Skip to content

Commit

Permalink
[SYCL] Fix CUDA tests using bfloat16 (intel#1421)
Browse files Browse the repository at this point in the history
* [SYCL] Fix CUDA tests using bfloat16
* Add missing using in element_wise_wi_marray_legacy

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
  • Loading branch information
steffenlarsen authored and bb-sycl committed Dec 7, 2022
1 parent 0aa5da0 commit f0fcbe5
Show file tree
Hide file tree
Showing 4 changed files with 3 additions and 5 deletions.
1 change: 1 addition & 0 deletions SYCL/BFloat16/bfloat16_builtins.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -13,6 +13,7 @@

using namespace sycl;
using namespace sycl::ext::oneapi;
using namespace sycl::ext::oneapi::experimental;

constexpr int N = 60; // divisible by all tested array sizes
constexpr float bf16_eps = 0.00390625;
Expand Down
1 change: 1 addition & 0 deletions SYCL/Matrix/element_wise_all_ops_bf16.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -17,6 +17,7 @@

using namespace sycl;
using namespace sycl::ext::intel;
using namespace sycl::ext::oneapi;
using namespace sycl::ext::oneapi::experimental::matrix;

#define SG_SZ 16
Expand Down
1 change: 1 addition & 0 deletions SYCL/Matrix/element_wise_wi_marray_legacy.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -13,6 +13,7 @@
#include <sycl/sycl.hpp>

using namespace sycl;
using namespace sycl::ext::oneapi;
using namespace sycl::ext::oneapi::experimental;
using namespace sycl::ext::oneapi::experimental::matrix;

Expand Down
5 changes: 0 additions & 5 deletions SYCL/Matrix/joint_matrix_tensorcores_legacy.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -66,13 +66,8 @@ T2 matrix_ref_mn(const int &m, const int &n, T1 *A, T1 *B, T2 *C) {
if constexpr (std::is_same<T1, uint16_t>::value) {
for (int k = 0; k < Big_K; k++)
res += make_fp32(A[m * Big_K + k]) * make_fp32(B[k * Big_N + n]);
} else if constexpr (std::is_same<T1, bfloat16>::value) {
for (int k = 0; k < Big_K; k++)
res +=
make_fp32(A[m * Big_K + k].raw()) * make_fp32(B[k * Big_N + n].raw());
} else {
for (int k = 0; k < Big_K; k++)

res +=
static_cast<T2>(A[m * Big_K + k]) * static_cast<T2>(B[k * Big_N + n]);
}
Expand Down

0 comments on commit f0fcbe5

Please sign in to comment.