From 90ea38c038095a7ab62eb2faf8d372347a7f2503 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Mon, 21 Mar 2022 18:10:13 +0300 Subject: [PATCH 1/8] [SYCL] Implement bf16 conversions on host device They are implemented in a way of RNE conversion. Signed-off-by: Dmitry Sidorov --- .../sycl/ext/intel/experimental/bfloat16.hpp | 19 +++- sycl/test/extensions/a.out | Bin 0 -> 14480 bytes sycl/test/extensions/bfloat16_host.cpp | 87 ++++++++++++++++++ 3 files changed, 102 insertions(+), 4 deletions(-) create mode 100755 sycl/test/extensions/a.out create mode 100644 sycl/test/extensions/bfloat16_host.cpp diff --git a/sycl/include/sycl/ext/intel/experimental/bfloat16.hpp b/sycl/include/sycl/ext/intel/experimental/bfloat16.hpp index 5a51f3746e225..388c858bc3cc4 100644 --- a/sycl/include/sycl/ext/intel/experimental/bfloat16.hpp +++ b/sycl/include/sycl/ext/intel/experimental/bfloat16.hpp @@ -31,16 +31,27 @@ class [[sycl_detail::uses_aspects(ext_intel_bf16_conversion)]] bfloat16 { #if defined(__SYCL_DEVICE_ONLY__) return __spirv_ConvertFToBF16INTEL(a); #else - throw exception{errc::feature_not_supported, - "Bfloat16 conversion is not supported on host device"}; + // In case if float value is nan - propagate bfloat16's qnan + if (std::isnan(a)) + return 0xffc1; + union { + uint32_t intStorage; + float floatValue; + }; + floatValue = a; + // Do RNE and truncate + uint32_t roundingBias = ((intStorage >> 16) & 0x1) + 0x00007FFF; + return static_cast((intStorage + roundingBias) >> 16); #endif } static float to_float(const storage_t &a) { #if defined(__SYCL_DEVICE_ONLY__) return __spirv_ConvertBF16ToFINTEL(a); #else - throw exception{errc::feature_not_supported, - "Bfloat16 conversion is not supported on host device"}; + // Shift temporary variable to silence the warning + uint32_t bits = a; + bits <<= 16; + return static_cast(bits); #endif } diff --git a/sycl/test/extensions/a.out b/sycl/test/extensions/a.out new file mode 100755 index 0000000000000000000000000000000000000000..7ff6910bbc1e22394c3224d0985ab0e58a2f43a9 GIT binary patch literal 14480 zcmeHOeQaCTb-yAlOa2HYD|TWh?fjg^h3&u;DN8bwAPXsp_ERH&)Ur~yoBK$Sj}jY- zWXMNH-jGz8o2-Q6)eF$VX#X(=v@kYoHPE6?vn{6Kpx&@QssQ^)v8)gVyP&%?Q)}1t z635!_+;eM|L=uKZ_8 z{Ya@hlOFHy?am~3X42Wx^v>x}|IYqie=+Cpk?kh?B)fP27>Nt&R%NQjuSSl?1v5;v zwAfw|z6QT7AKzTH{><9J?a_Z2kxET2zXI0P1poFDco!;cCjZkVy|K?Fiu4)nRx6m8k>G$ z7!y;utWmV$18!6$~aO;|0Si#M4%BI1!B=0=5el(g`D1v7n#0n)N|7K2%Owdxi^PZGRnGg-1N4%I7!7$c)@;cn=}|wgZR3TH^7iy^z)lr;Tb} zv;m-X#7T4Ds!Cb3&4D}DXTX7b*$8OJf%D!&#bF2T+@Ey^?%baaI&kup(-8;0mO)Vc z#QW{~$=2v5P3W&wEYHHMe&X%+%VME#0MNqLAzV8)hmj+Gk~FW+Eg)@uf_Q53^;yaH z5Km3LUXlC`;=RP5m;58dQ`4@WmHb1*Q$P^-B;rdbwlsSEL5%S1%}N2Z6qW&9uBM(RPr{uMK&3f<+a)Y}3zM~fjT$q!3M=z%&eT%R${pAmof_~|)S1-@& zm#R4|^d&UoN3FlvsEHQp0_h)lRris^j{aL2 zmDl?Rb0csidWRe>zj#|Om*#JDbhiE}#`s3Zy69E?)QeU9ROxDUmRbU)+J@<%U*sj! zabdpNcNgPQzM{Xy`)$=r2nl~kS~cP7Uw-VC{0R`MuuoO^9@OSg;d@oXu5hmUjRqT3 z;oFPMyiN!-tK|ma>Sq>7P67#5F#iKbFU$FL;(bk5E!R&*(F4({e)7e6S%0d030O5% zZij+~Zg&hf`d8|cF7!Gs(6BTqKkP`btrXllKI$@5a#}Ap$?~50oW1Q*e=Lbi+lTrwGGOb`fuT43{WchPuA!}x| ze7q2!(sFr=l|j#9TW}`6H!!VmId&aNnxb=zOSFFR$frLcUQ^|-yJ%g~&dn-&+thRJ zLB%uW`k1%xqSyzOoZFt)^C|4bEy9O8DfKL+b;zwyl%^!{w1Ka|n*_DcAL*Tam3srvW2)XGz$vvn?0u4+#p%92^B)TfRZ#82g#~&pixGvPJZJ!P z29%zoPlKL=Tm|&+Ne|O10Qs3a3kwH9E1-GMIndLf0_*)eXdZMH^fc%^XazI|djV`! z=Rhf~w~1)^y+fj9y0hiswQcRETiQCwKYEv%gO9WfDK(e1+d{HeA*;-aN5Kd1JAHd$ z;Vh`|b?)(P{7lE17usjU(+@qlW7ngP5KMa1H|@|35fJfpzUmpgZ)FjVfTi+j{B*Qk zc0QMTK9T}deg?m7=o9(M$>o{4{u`u^@^zp!`y=)8cUAer&hlsL<+oLN%31D-*XrMh zzC8o`OX-JH`BtZXtZx4Z%60S;?R+)=_Sogr=hG+;pQJ;2d#ei?yYpuF3@ajt47dp|l&lfKELJ?me;@dXh(;x@= z0obSUr#VA&jmMPCEvMxkSnh%49$4;yf#n`p?t$eVSnh%49$4;y|KB}uzAY>r zJdg8SN@o>cQFKP(c;%DHO@HXZaV*s!_l}FFZyS{OO#a-;u;kTeN@1yUo=p;^@2oMI zlkd>TZB+O}iq5KXdf%r+@28Z0eRm;8blMXpB3k7Y=N8_u%<(sba~_-LZydwQKl#oL zNv>~$)T&>+oW-#W1bQ}mXiE7kmZK+#Vt+OOyyMW0pFZFleB;1k-m zvGG#YDrvj@eg44CT_wo{UkZl&f!^&lDiWD^c0xO57K-UyRtrMVAJDcmFN_a=W8S0TASC6_)5`iK6pi&;Ca+2 z-!6DQHsY%U&)Y_Pwcz#Jh_4Y9HO`IrS`lk9o?dZZZQVD@`^1de4;t}~+WBcC-YEj= z{I?NbSBrmW#Jj|7Q~CRA`%a_$dcpfmBfbG~9Wf^)^Iq|Q;P0`{La)GXp!SVMe4~gp z#p~bc8|8TytfRDjtq{Lls3UfcBt6)5x&K2-fX-rD5C^~?5Y_f4o>uZ_)p+rKI03ww z{g;$H-%ofye_h)5iSz2bTLVw&FYqIOD#|{~zYUyD-TB`*^anvr;Yejc%5aYo6z zzc+kO;`MpU_piTm$;XtPp962!|Gx%K{n@GR*GgRcM&V%>z8W_UlIQz7%WJ@CKAhfS z*9di)x+T6|yz83h&ni3o{Gi0exRk$NR9x~)`Tw$%?-F-9!fc{Oo3;0Akf$#&9M8nU zeZ9cjBplK2tl>cUK6Z0e7$8_aVOy^NRN| zh{^O3@GhGFJTBbP-sbT-I`DO#jeJ6_BOaJ7z7v#F( z58;}^`FU|%;X>iO4}1f7v-qb!2TtSb9>+J8ompjPtE^p!3vxbe5@)u=6k?N`($1zu z{tJu?`FY;u=UO>l4JITLlsLvk?LVdxi7X+11UT8RtP0!uO8%s>Ls1-*cznOLgq=5m zYf`fD^0!K!Lr-d@L1B*-W@w-S*jrnxOT(TSyjnAPl zU&v7}INdP_Gm55FG~$JVppWauK|rIXwLd;(E&>p*_Fox5tmLBB2K>pm6?dzdS*wsX z?OL0tq_e5qVpDclfRUbxPnboKC|E_SluG#%B5BrQnt$3+T(2BwiLO-&t#ii4Ars6o^)_8f}riyG1W zLllo_AYjyp>I&3{4hduE^ZO(Fh6hnbQ?90*Yy@NopD@&XJg{faXcQrck%5sYteSS8 z4j(|IUoxFFN<}j%WVqk|JgzYuc7_L{_lh}VGL8sI6OPGA;|f8fSWAkjOu00Z6u_9t zB@mKGfru7@DD!DDQjb$q8Y(PNnXio@0zYbjlN_+=s>3+EZ;-qA3;TN#nO^d6mpN_q z$@iyVk2#HXnV!P<#WNvZ$H73Vkef2(77;!4jEv8G=AN|^$5M7E>AkD$D#TxqIw~~O z-3%H>J6e-ke9xYQtX!=T6!bY>EY_@yq6>=Ic-9U+bp^vxpsL#kba%o?6id7z+M{Z@ z#Qq!J4ffc5YsaFJWm-^hQ`S_EO;t72X2HEg@-})!d&HC#9|yGxHl5@g&4WT-__I07 z^hXAUcUp05#uM3+e==U26#nG#ER=0(73>oAdfvboD8qb?lY+|SGnVkn`-LB2x5&%( z`3pI<)eVkB_U|>YW=9T0hKzx+{XzIH?W-(~3eg&qm{KF#L`{C^I7-o>@><0JE* zLWbU*Sf9@s`2Pqr7E~0ED-9_K?iS9JSby3RmgUn5;P)NYYnP?<%ah2`J0t6#^Mt`6 z21NlB`$7qi6SAGG&+lD>!11V8)1UK(!O{7hB!bWBY|gg9d7{tYM`vVQKc9o_pbC-r z&6)jYIi~d9MrUr!^LdG`^uuHXiCB_h{|TdHzbDuI=W~^>Dt+A+!@HcFSfAVfJYXst zQ0097GOP68RXEAe*`izj_kmFt*UmXthQXo#6~O0_62yi5VEz~~PW_6~uPFVqET|~I zw=#X%r5|exOUameq3?24$a>10_(PXIp9@&fmUr`)YWh;W_MXfC&-0oM*RJmN z|JbF^=UlU@eu^V>CZ0cx{R9~wermbZ&cpcsdw8C*JnJ*PhJp@U*rFuqIgJ3YB|*AKimBfWz_$y&*yOSn@OwwoxtYm>xb*712m_8?YvI2 zHDWfYU0JV2q3qP>_sN)MtFWJJfc2R7L76c2e@2VRqP!-F*ht15Nps{Wdq2uaB;dkX z|D+3Nb581-BDTgBrKlDl-G}I`kp1U*LB;I5(y$4=@FRV|181|+SCy4y9j?vvKj^ll MKEaMEB^N9H6VH +#include + +#include +#include +#include +#include +#include + +using sycl::ext::intel::experimental::bfloat16; + +// Helper to convert the expected bits to float value to compare with the result +typedef union { + float Value; + struct { + uint32_t Mantissa : 23; + uint32_t Exponent : 8; + uint32_t Sign : 1; + } RawData; +} floatConvHelper; + +float bistToFloatConv(std::string &Bits) { + floatConvHelper &Helper; + Helper.RawData.Sign = static_cast(Bits[0] - '0'); + uint32_t Exponent = 0; + for (size_t I = 1; I != 9; ++I) + Exponent = Exponent + static_cast(Bits[I] - '0') * pow(2, 8 - I); + Helper.RawData.Exponent = Exponent; + uint32_t Mantissa = 0; + for (size_t I = 9; I != 32; ++I) + Mantissa = Mantissa + static_cast(Bits[I] - '0') * pow(2, 8 - I); + Helper.RawData.Mantissa = Mantissa; +} + +inline bool check_bf16_from_float(float &Val, uint16_t &Expected) { + if (from_float(Val) != Expected) { + std::cout << "from_float check for Val = " << Val << " failed!\n"; + return false; + } + return true; +} + +inline bool check_bf16_to_float(uint16_t &Val, float &Expected) { + if (to_float(Val) != Expected) { + std::cout << "to_float check for Val = " << Val << " failed!\n"; + return false; + } + return true; +} + +int main() { + bool Success = + check_bf16_from_float(0.0f, std::stoi("0000000000000000", nullptr, 2)); + Success &= check_bf16_from_float(42.0f, + std::stoi("100001000101000", nullptr, 2)); + Success &= check_bf16_from_float(std::numeric_limits::min(), + std::stoi("0000000010000000", nullptr, 2)); + Success &= check_bf16_from_float(std::numeric_limits::max(), + std::stoi("0111111110000000", nullptr, 2)); + Success &= check_bf16_from_float(std::numeric_limits::quiet_NaN(), + std::stoi("1111111111000001", nullptr, 2)); + + Success &= + check_bf16_to_float(to_float(0), + bitToFloatConv("00000000000000000000000000000000")); + Success &= + check_bf16_to_float(to_float(1), + bitToFloatConv("01000111100000000000000000000000")); + Success &= + check_bf16_to_float(to_float(42), + bitToFloatConv("00000000001010100000000000000000")); + Success &= + check_bf16_to_float(to_float(std::numeric_limits::max()), + bitToFloatConv("11111111111111110000000000000000")); + if (!Success) + return -1; + return 0; +} From a91a5b53f5a6ae07a99e87c20e12677cfff998ec Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Sun, 3 Apr 2022 22:02:51 +0300 Subject: [PATCH 2/8] Clang-format Signed-off-by: Dmitry Sidorov --- sycl/test/extensions/bfloat16_host.cpp | 21 +++++++++------------ 1 file changed, 9 insertions(+), 12 deletions(-) diff --git a/sycl/test/extensions/bfloat16_host.cpp b/sycl/test/extensions/bfloat16_host.cpp index fc50e37c602ee..01c698fde2654 100644 --- a/sycl/test/extensions/bfloat16_host.cpp +++ b/sycl/test/extensions/bfloat16_host.cpp @@ -10,8 +10,8 @@ #include #include -#include #include +#include #include #include #include @@ -60,8 +60,8 @@ inline bool check_bf16_to_float(uint16_t &Val, float &Expected) { int main() { bool Success = check_bf16_from_float(0.0f, std::stoi("0000000000000000", nullptr, 2)); - Success &= check_bf16_from_float(42.0f, - std::stoi("100001000101000", nullptr, 2)); + Success &= + check_bf16_from_float(42.0f, std::stoi("100001000101000", nullptr, 2)); Success &= check_bf16_from_float(std::numeric_limits::min(), std::stoi("0000000010000000", nullptr, 2)); Success &= check_bf16_from_float(std::numeric_limits::max(), @@ -69,15 +69,12 @@ int main() { Success &= check_bf16_from_float(std::numeric_limits::quiet_NaN(), std::stoi("1111111111000001", nullptr, 2)); - Success &= - check_bf16_to_float(to_float(0), - bitToFloatConv("00000000000000000000000000000000")); - Success &= - check_bf16_to_float(to_float(1), - bitToFloatConv("01000111100000000000000000000000")); - Success &= - check_bf16_to_float(to_float(42), - bitToFloatConv("00000000001010100000000000000000")); + Success &= check_bf16_to_float( + to_float(0), bitToFloatConv("00000000000000000000000000000000")); + Success &= check_bf16_to_float( + to_float(1), bitToFloatConv("01000111100000000000000000000000")); + Success &= check_bf16_to_float( + to_float(42), bitToFloatConv("00000000001010100000000000000000")); Success &= check_bf16_to_float(to_float(std::numeric_limits::max()), bitToFloatConv("11111111111111110000000000000000")); From 0636c0b6e05df8ee8fbbe62e81f64dc0e3938cff Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Mon, 4 Apr 2022 20:13:17 +0300 Subject: [PATCH 3/8] Fix test Signed-off-by: Dmitry Sidorov --- sycl/test/extensions/bfloat16_host.cpp | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/sycl/test/extensions/bfloat16_host.cpp b/sycl/test/extensions/bfloat16_host.cpp index 01c698fde2654..7fec5f5748b31 100644 --- a/sycl/test/extensions/bfloat16_host.cpp +++ b/sycl/test/extensions/bfloat16_host.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// +// RUN: %clangxx -fsycl %s -o %t.out // RUN: %RUN_ON_HOST %t.out #include #include @@ -28,7 +29,7 @@ typedef union { } RawData; } floatConvHelper; -float bistToFloatConv(std::string &Bits) { +float bitsToFloatConv(std::string &Bits) { floatConvHelper &Helper; Helper.RawData.Sign = static_cast(Bits[0] - '0'); uint32_t Exponent = 0; @@ -70,14 +71,14 @@ int main() { std::stoi("1111111111000001", nullptr, 2)); Success &= check_bf16_to_float( - to_float(0), bitToFloatConv("00000000000000000000000000000000")); + to_float(0), bitsToFloatConv("00000000000000000000000000000000")); Success &= check_bf16_to_float( - to_float(1), bitToFloatConv("01000111100000000000000000000000")); + to_float(1), bitsToFloatConv("01000111100000000000000000000000")); Success &= check_bf16_to_float( - to_float(42), bitToFloatConv("00000000001010100000000000000000")); + to_float(42), bitsToFloatConv("00000000001010100000000000000000")); Success &= check_bf16_to_float(to_float(std::numeric_limits::max()), - bitToFloatConv("11111111111111110000000000000000")); + bitsToFloatConv("11111111111111110000000000000000")); if (!Success) return -1; return 0; From 8055dd1f75fcfae2a14fc4a87229e35b148311e3 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Mon, 4 Apr 2022 22:09:55 +0300 Subject: [PATCH 4/8] Final fixes Signed-off-by: Dmitry Sidorov --- .../sycl/ext/intel/experimental/bfloat16.hpp | 2 + sycl/test/extensions/a.out | Bin 14480 -> 0 bytes sycl/test/extensions/bfloat16_host.cpp | 37 ++++++++++-------- 3 files changed, 22 insertions(+), 17 deletions(-) delete mode 100755 sycl/test/extensions/a.out diff --git a/sycl/include/sycl/ext/intel/experimental/bfloat16.hpp b/sycl/include/sycl/ext/intel/experimental/bfloat16.hpp index 388c858bc3cc4..c3c51b452f1d1 100644 --- a/sycl/include/sycl/ext/intel/experimental/bfloat16.hpp +++ b/sycl/include/sycl/ext/intel/experimental/bfloat16.hpp @@ -11,6 +11,8 @@ #include #include +#include + __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { namespace ext { diff --git a/sycl/test/extensions/a.out b/sycl/test/extensions/a.out deleted file mode 100755 index 7ff6910bbc1e22394c3224d0985ab0e58a2f43a9..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 14480 zcmeHOeQaCTb-yAlOa2HYD|TWh?fjg^h3&u;DN8bwAPXsp_ERH&)Ur~yoBK$Sj}jY- zWXMNH-jGz8o2-Q6)eF$VX#X(=v@kYoHPE6?vn{6Kpx&@QssQ^)v8)gVyP&%?Q)}1t z635!_+;eM|L=uKZ_8 z{Ya@hlOFHy?am~3X42Wx^v>x}|IYqie=+Cpk?kh?B)fP27>Nt&R%NQjuSSl?1v5;v zwAfw|z6QT7AKzTH{><9J?a_Z2kxET2zXI0P1poFDco!;cCjZkVy|K?Fiu4)nRx6m8k>G$ z7!y;utWmV$18!6$~aO;|0Si#M4%BI1!B=0=5el(g`D1v7n#0n)N|7K2%Owdxi^PZGRnGg-1N4%I7!7$c)@;cn=}|wgZR3TH^7iy^z)lr;Tb} zv;m-X#7T4Ds!Cb3&4D}DXTX7b*$8OJf%D!&#bF2T+@Ey^?%baaI&kup(-8;0mO)Vc z#QW{~$=2v5P3W&wEYHHMe&X%+%VME#0MNqLAzV8)hmj+Gk~FW+Eg)@uf_Q53^;yaH z5Km3LUXlC`;=RP5m;58dQ`4@WmHb1*Q$P^-B;rdbwlsSEL5%S1%}N2Z6qW&9uBM(RPr{uMK&3f<+a)Y}3zM~fjT$q!3M=z%&eT%R${pAmof_~|)S1-@& zm#R4|^d&UoN3FlvsEHQp0_h)lRris^j{aL2 zmDl?Rb0csidWRe>zj#|Om*#JDbhiE}#`s3Zy69E?)QeU9ROxDUmRbU)+J@<%U*sj! zabdpNcNgPQzM{Xy`)$=r2nl~kS~cP7Uw-VC{0R`MuuoO^9@OSg;d@oXu5hmUjRqT3 z;oFPMyiN!-tK|ma>Sq>7P67#5F#iKbFU$FL;(bk5E!R&*(F4({e)7e6S%0d030O5% zZij+~Zg&hf`d8|cF7!Gs(6BTqKkP`btrXllKI$@5a#}Ap$?~50oW1Q*e=Lbi+lTrwGGOb`fuT43{WchPuA!}x| ze7q2!(sFr=l|j#9TW}`6H!!VmId&aNnxb=zOSFFR$frLcUQ^|-yJ%g~&dn-&+thRJ zLB%uW`k1%xqSyzOoZFt)^C|4bEy9O8DfKL+b;zwyl%^!{w1Ka|n*_DcAL*Tam3srvW2)XGz$vvn?0u4+#p%92^B)TfRZ#82g#~&pixGvPJZJ!P z29%zoPlKL=Tm|&+Ne|O10Qs3a3kwH9E1-GMIndLf0_*)eXdZMH^fc%^XazI|djV`! z=Rhf~w~1)^y+fj9y0hiswQcRETiQCwKYEv%gO9WfDK(e1+d{HeA*;-aN5Kd1JAHd$ z;Vh`|b?)(P{7lE17usjU(+@qlW7ngP5KMa1H|@|35fJfpzUmpgZ)FjVfTi+j{B*Qk zc0QMTK9T}deg?m7=o9(M$>o{4{u`u^@^zp!`y=)8cUAer&hlsL<+oLN%31D-*XrMh zzC8o`OX-JH`BtZXtZx4Z%60S;?R+)=_Sogr=hG+;pQJ;2d#ei?yYpuF3@ajt47dp|l&lfKELJ?me;@dXh(;x@= z0obSUr#VA&jmMPCEvMxkSnh%49$4;yf#n`p?t$eVSnh%49$4;y|KB}uzAY>r zJdg8SN@o>cQFKP(c;%DHO@HXZaV*s!_l}FFZyS{OO#a-;u;kTeN@1yUo=p;^@2oMI zlkd>TZB+O}iq5KXdf%r+@28Z0eRm;8blMXpB3k7Y=N8_u%<(sba~_-LZydwQKl#oL zNv>~$)T&>+oW-#W1bQ}mXiE7kmZK+#Vt+OOyyMW0pFZFleB;1k-m zvGG#YDrvj@eg44CT_wo{UkZl&f!^&lDiWD^c0xO57K-UyRtrMVAJDcmFN_a=W8S0TASC6_)5`iK6pi&;Ca+2 z-!6DQHsY%U&)Y_Pwcz#Jh_4Y9HO`IrS`lk9o?dZZZQVD@`^1de4;t}~+WBcC-YEj= z{I?NbSBrmW#Jj|7Q~CRA`%a_$dcpfmBfbG~9Wf^)^Iq|Q;P0`{La)GXp!SVMe4~gp z#p~bc8|8TytfRDjtq{Lls3UfcBt6)5x&K2-fX-rD5C^~?5Y_f4o>uZ_)p+rKI03ww z{g;$H-%ofye_h)5iSz2bTLVw&FYqIOD#|{~zYUyD-TB`*^anvr;Yejc%5aYo6z zzc+kO;`MpU_piTm$;XtPp962!|Gx%K{n@GR*GgRcM&V%>z8W_UlIQz7%WJ@CKAhfS z*9di)x+T6|yz83h&ni3o{Gi0exRk$NR9x~)`Tw$%?-F-9!fc{Oo3;0Akf$#&9M8nU zeZ9cjBplK2tl>cUK6Z0e7$8_aVOy^NRN| zh{^O3@GhGFJTBbP-sbT-I`DO#jeJ6_BOaJ7z7v#F( z58;}^`FU|%;X>iO4}1f7v-qb!2TtSb9>+J8ompjPtE^p!3vxbe5@)u=6k?N`($1zu z{tJu?`FY;u=UO>l4JITLlsLvk?LVdxi7X+11UT8RtP0!uO8%s>Ls1-*cznOLgq=5m zYf`fD^0!K!Lr-d@L1B*-W@w-S*jrnxOT(TSyjnAPl zU&v7}INdP_Gm55FG~$JVppWauK|rIXwLd;(E&>p*_Fox5tmLBB2K>pm6?dzdS*wsX z?OL0tq_e5qVpDclfRUbxPnboKC|E_SluG#%B5BrQnt$3+T(2BwiLO-&t#ii4Ars6o^)_8f}riyG1W zLllo_AYjyp>I&3{4hduE^ZO(Fh6hnbQ?90*Yy@NopD@&XJg{faXcQrck%5sYteSS8 z4j(|IUoxFFN<}j%WVqk|JgzYuc7_L{_lh}VGL8sI6OPGA;|f8fSWAkjOu00Z6u_9t zB@mKGfru7@DD!DDQjb$q8Y(PNnXio@0zYbjlN_+=s>3+EZ;-qA3;TN#nO^d6mpN_q z$@iyVk2#HXnV!P<#WNvZ$H73Vkef2(77;!4jEv8G=AN|^$5M7E>AkD$D#TxqIw~~O z-3%H>J6e-ke9xYQtX!=T6!bY>EY_@yq6>=Ic-9U+bp^vxpsL#kba%o?6id7z+M{Z@ z#Qq!J4ffc5YsaFJWm-^hQ`S_EO;t72X2HEg@-})!d&HC#9|yGxHl5@g&4WT-__I07 z^hXAUcUp05#uM3+e==U26#nG#ER=0(73>oAdfvboD8qb?lY+|SGnVkn`-LB2x5&%( z`3pI<)eVkB_U|>YW=9T0hKzx+{XzIH?W-(~3eg&qm{KF#L`{C^I7-o>@><0JE* zLWbU*Sf9@s`2Pqr7E~0ED-9_K?iS9JSby3RmgUn5;P)NYYnP?<%ah2`J0t6#^Mt`6 z21NlB`$7qi6SAGG&+lD>!11V8)1UK(!O{7hB!bWBY|gg9d7{tYM`vVQKc9o_pbC-r z&6)jYIi~d9MrUr!^LdG`^uuHXiCB_h{|TdHzbDuI=W~^>Dt+A+!@HcFSfAVfJYXst zQ0097GOP68RXEAe*`izj_kmFt*UmXthQXo#6~O0_62yi5VEz~~PW_6~uPFVqET|~I zw=#X%r5|exOUameq3?24$a>10_(PXIp9@&fmUr`)YWh;W_MXfC&-0oM*RJmN z|JbF^=UlU@eu^V>CZ0cx{R9~wermbZ&cpcsdw8C*JnJ*PhJp@U*rFuqIgJ3YB|*AKimBfWz_$y&*yOSn@OwwoxtYm>xb*712m_8?YvI2 zHDWfYU0JV2q3qP>_sN)MtFWJJfc2R7L76c2e@2VRqP!-F*ht15Nps{Wdq2uaB;dkX z|D+3Nb581-BDTgBrKlDl-G}I`kp1U*LB;I5(y$4=@FRV|181|+SCy4y9j?vvKj^ll MKEaMEB^N9H6VH #include -using sycl::ext::intel::experimental::bfloat16; - // Helper to convert the expected bits to float value to compare with the result typedef union { float Value; @@ -29,8 +27,8 @@ typedef union { } RawData; } floatConvHelper; -float bitsToFloatConv(std::string &Bits) { - floatConvHelper &Helper; +float bitsToFloatConv(std::string Bits) { + floatConvHelper Helper; Helper.RawData.Sign = static_cast(Bits[0] - '0'); uint32_t Exponent = 0; for (size_t I = 1; I != 9; ++I) @@ -38,21 +36,26 @@ float bitsToFloatConv(std::string &Bits) { Helper.RawData.Exponent = Exponent; uint32_t Mantissa = 0; for (size_t I = 9; I != 32; ++I) - Mantissa = Mantissa + static_cast(Bits[I] - '0') * pow(2, 8 - I); + Mantissa = Mantissa + static_cast(Bits[I] - '0') * pow(2, 31 - I); Helper.RawData.Mantissa = Mantissa; + return Helper.Value; } -inline bool check_bf16_from_float(float &Val, uint16_t &Expected) { - if (from_float(Val) != Expected) { - std::cout << "from_float check for Val = " << Val << " failed!\n"; +bool check_bf16_from_float(float Val, uint16_t Expected) { + uint16_t Result = sycl::ext::intel::experimental::bfloat16::from_float(Val); + if (Result != Expected) { + std::cout << "from_float check for Val = " << Val << " failed!\n" + << "Expected " << Expected << " Got " << Result << "\n"; return false; } return true; } -inline bool check_bf16_to_float(uint16_t &Val, float &Expected) { - if (to_float(Val) != Expected) { - std::cout << "to_float check for Val = " << Val << " failed!\n"; +bool check_bf16_to_float(uint16_t Val, float Expected) { + float Result = sycl::ext::intel::experimental::bfloat16::to_float(Val); + if (Result != Expected) { + std::cout << "to_float check for Val = " << Val << " failed!\n" + << "Expected " << Expected << " Got " << Result << "\n"; return false; } return true; @@ -71,14 +74,14 @@ int main() { std::stoi("1111111111000001", nullptr, 2)); Success &= check_bf16_to_float( - to_float(0), bitsToFloatConv("00000000000000000000000000000000")); + 0, bitsToFloatConv(std::string("00000000000000000000000000000000"))); Success &= check_bf16_to_float( - to_float(1), bitsToFloatConv("01000111100000000000000000000000")); + 1, bitsToFloatConv(std::string("01000111100000000000000000000000"))); Success &= check_bf16_to_float( - to_float(42), bitsToFloatConv("00000000001010100000000000000000")); - Success &= - check_bf16_to_float(to_float(std::numeric_limits::max()), - bitsToFloatConv("11111111111111110000000000000000")); + 42, bitsToFloatConv(std::string("01001010001010000000000000000000"))); + Success &= check_bf16_to_float( + std::numeric_limits::max(), + bitsToFloatConv(std::string("01001111011111111111111100000000"))); if (!Success) return -1; return 0; From 1c4ccccb8229568916f53cd36973706a2c0f8c04 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Tue, 5 Apr 2022 22:59:34 +0300 Subject: [PATCH 5/8] Apply intel/llvm review change Signed-off-by: Dmitry Sidorov --- sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp b/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp index c3dad5eefe0ed..78fe2c3155abe 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp @@ -11,7 +11,9 @@ #include #include +#if !defined(__SYCL_DEVICE_ONLY__) #include +#endif __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { From e5c66f4fe922a3abc0984f71f5e0e58a783551e1 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Fri, 29 Apr 2022 04:57:39 -0700 Subject: [PATCH 6/8] Fix extension location and namespace in tests Signed-off-by: Larsen, Steffen --- sycl/test/extensions/bfloat16_host.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/test/extensions/bfloat16_host.cpp b/sycl/test/extensions/bfloat16_host.cpp index d4498e746dfb2..e3cfb71abb558 100644 --- a/sycl/test/extensions/bfloat16_host.cpp +++ b/sycl/test/extensions/bfloat16_host.cpp @@ -8,7 +8,7 @@ // RUN: %clangxx -fsycl %s -o %t.out // RUN: %RUN_ON_HOST %t.out -#include +#include #include #include @@ -42,7 +42,7 @@ float bitsToFloatConv(std::string Bits) { } bool check_bf16_from_float(float Val, uint16_t Expected) { - uint16_t Result = sycl::ext::intel::experimental::bfloat16::from_float(Val); + uint16_t Result = sycl::ext::oneapi::experimental::bfloat16::from_float(Val); if (Result != Expected) { std::cout << "from_float check for Val = " << Val << " failed!\n" << "Expected " << Expected << " Got " << Result << "\n"; @@ -52,7 +52,7 @@ bool check_bf16_from_float(float Val, uint16_t Expected) { } bool check_bf16_to_float(uint16_t Val, float Expected) { - float Result = sycl::ext::intel::experimental::bfloat16::to_float(Val); + float Result = sycl::ext::oneapi::experimental::bfloat16::to_float(Val); if (Result != Expected) { std::cout << "to_float check for Val = " << Val << " failed!\n" << "Expected " << Expected << " Got " << Result << "\n"; From 81c7411ba198a1c508719618e55bab102eddddee Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Mon, 22 Aug 2022 10:25:49 -0700 Subject: [PATCH 7/8] Fix namespace Signed-off-by: Sidorov, Dmitry --- sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp b/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp index b1aefff079100..42a95fac59bf4 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp @@ -16,6 +16,7 @@ #endif namespace sycl { +__SYCL_INLINE_VER_NAMESPACE(_V1) { namespace ext { namespace oneapi { namespace experimental { From cf7d68da93695a87283b3266d9ee55008d1ed6fd Mon Sep 17 00:00:00 2001 From: "Sidorov, Dmitry" Date: Mon, 22 Aug 2022 11:51:46 -0700 Subject: [PATCH 8/8] Grammar fix Signed-off-by: Sidorov, Dmitry --- sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp b/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp index 42a95fac59bf4..3c97bda5b4e90 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/bfloat16.hpp @@ -39,7 +39,7 @@ class bfloat16 { return __spirv_ConvertFToBF16INTEL(a); #endif #else - // In case if float value is nan - propagate bfloat16's qnan + // In case of float value is nan - propagate bfloat16's qnan if (std::isnan(a)) return 0xffc1; union {