|
| 1 | + |
| 2 | +#define TM 8 |
| 3 | +#define TN SG_SZ |
| 4 | +#define TK 16 |
| 5 | + |
| 6 | +static float make_fp32(bfloat16 x) { |
| 7 | + unsigned int y = sycl::bit_cast<uint16_t>(x); |
| 8 | + y = y << 16; |
| 9 | + float *res = reinterpret_cast<float *>(&y); |
| 10 | + return *res; |
| 11 | +} |
| 12 | + |
| 13 | +template <typename T, size_t NUM_ROWS, size_t NUM_COLS> struct big_matrix { |
| 14 | +public: |
| 15 | + T *mat; |
| 16 | + |
| 17 | +public: |
| 18 | + T *get_data() { return mat; } |
| 19 | + void set_data(T *data) { mat = data; } |
| 20 | + big_matrix(T *data) : mat(data) {} |
| 21 | +}; |
| 22 | + |
| 23 | +template <typename T> struct apply_add { |
| 24 | + void operator()(T &x) const { x = x + bfloat16(2); } |
| 25 | +}; |
| 26 | + |
| 27 | +template <typename T, size_t M, size_t N, typename F> |
| 28 | +void matrix_verify_add(queue q, big_matrix<T, M, N> &A, nd_range<2> &r, |
| 29 | + const float ref, F &&lambda) { |
| 30 | + buffer<bfloat16, 2> bufA(A.get_data(), range<2>(M, N)); |
| 31 | + |
| 32 | + q.submit([&](handler &cgh) { |
| 33 | + accessor accA{bufA, cgh}; |
| 34 | + |
| 35 | + cgh.parallel_for(r, [accA, lambda]( |
| 36 | + nd_item<2> spmd_item) [[intel::reqd_sub_group_size( |
| 37 | + SG_SZ)]] { |
| 38 | + const auto global_idx = spmd_item.get_global_id(0); |
| 39 | + const auto global_idy = spmd_item.get_global_id(1); |
| 40 | + const auto sg_startx = global_idx - spmd_item.get_local_id(0); |
| 41 | + const auto sg_starty = global_idy - spmd_item.get_local_id(1); |
| 42 | + |
| 43 | + sub_group sg = spmd_item.get_sub_group(); |
| 44 | + joint_matrix<sub_group, T, use::a, TM, TK, layout::row_major> sub_a; |
| 45 | + |
| 46 | + joint_matrix_fill(sg, sub_a, bfloat16(5.0)); |
| 47 | + |
| 48 | + joint_matrix_apply(sg, sub_a, lambda); |
| 49 | + |
| 50 | + ext::intel::experimental::matrix::joint_matrix_store( |
| 51 | + sg, sub_a, |
| 52 | + accA.get_pointer() + (sg_startx * TM) * N + sg_starty / SG_SZ * TN, |
| 53 | + N); |
| 54 | + }); // parallel for |
| 55 | + }).wait(); |
| 56 | + // Check if the results are correct |
| 57 | + { |
| 58 | + host_accessor Acc{bufA}; |
| 59 | + assert(std::all_of(Acc.begin(), Acc.end(), [=](auto Elem) { |
| 60 | + return (std::fabs(static_cast<float>(make_fp32(Elem) - ref)) < |
| 61 | + std::numeric_limits<float>::epsilon()); |
| 62 | + })); |
| 63 | + } |
| 64 | +} |
| 65 | + |
| 66 | +static constexpr size_t MATRIX_M = TM * 2; |
| 67 | +static constexpr size_t MATRIX_N = TN * 2; |
| 68 | +bfloat16 A[MATRIX_M][MATRIX_N]; |
| 69 | +float D[MATRIX_M][MATRIX_N]; |
| 70 | + |
| 71 | +void matrix_ops_ref(float *D, int M, int N) { |
| 72 | + for (int m = 0; m < M; m++) |
| 73 | + for (int n = 0; n < N; n++) { |
| 74 | + *(D + m * N + n) = 0; |
| 75 | + *(D + m * N + n) *= 2; |
| 76 | + } |
| 77 | +} |
| 78 | + |
| 79 | +int main() { |
| 80 | + |
| 81 | + big_matrix<float, MATRIX_M, MATRIX_N> MD((float *)&D); |
| 82 | + big_matrix<bfloat16, MATRIX_M, MATRIX_N> MA((bfloat16 *)&A); |
| 83 | + |
| 84 | + size_t NDRangeM = MATRIX_M / TM; |
| 85 | + size_t NDRangeN = MATRIX_N / TN; |
| 86 | + queue q; |
| 87 | + nd_range<2> r({NDRangeM, NDRangeN * SG_SZ}, {1, 1 * SG_SZ}); |
| 88 | + |
| 89 | + matrix_verify_add<bfloat16, MATRIX_M, MATRIX_N>( |
| 90 | + q, MA, r, 7.0, [=](bfloat16 &x) { x = x + bfloat16(2); }); |
| 91 | + matrix_verify_add<bfloat16, MATRIX_M, MATRIX_N>(q, MA, r, 7.0, |
| 92 | + apply_add<bfloat16>()); |
| 93 | + std::cout << "Passed\n"; |
| 94 | + return 0; |
| 95 | +} |
0 commit comments