This repository was archived by the owner on Mar 28, 2023. It is now read-only.
forked from llvm/llvm-test-suite
-
Notifications
You must be signed in to change notification settings - Fork 130
[SYCL][Matrix] Adding test cases for the joint_matrix_apply() and fixing namespace for get_wi_data() #1636
Merged
Merged
[SYCL][Matrix] Adding test cases for the joint_matrix_apply() and fixing namespace for get_wi_data() #1636
Changes from all commits
Commits
Show all changes
20 commits
Select commit
Hold shift + click to select a range
930bb3a
[SYCL][Matrix] Adding test cases for the joint_matrix_apply() interface
arnamoy10 8b7deaa
Changing file names and addressing reviewer comments.
arnamoy10 c3d0f2b
clang-format
arnamoy10 a91b20e
Using only one function for both lambda and function object.
arnamoy10 88f54e5
clang-format
arnamoy10 7a7f8d4
Inlining the comparison function.
arnamoy10 b886d04
clang-format
arnamoy10 aa04af9
Chanign the namespace for get_wi_data() use (using intel namespace
arnamoy10 a80c6fb
clang-format.
arnamoy10 97e9a05
Reverting legacy changes
arnamoy10 74d6a3a
clang-format
arnamoy10 0196b38
Reverting cuda changes
arnamoy10 3af272d
clang-format
arnamoy10 cb0d4fc
Reverting more CUDA changes
arnamoy10 253ccd7
clang-format
arnamoy10 e97508d
Address reviewers comments.
arnamoy10 ecbb1b8
clang-format
arnamoy10 5fd1cd2
Adding sycl::bit_cast
arnamoy10 7f17022
clang-format
arnamoy10 3812e97
clang-format
arnamoy10 File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,24 @@ | ||
| //==----------- joint_matrix_apply_bf16.cpp - DPC++ joint_matrix-----------==// | ||
| // | ||
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
| // See https://llvm.org/LICENSE.txt for license information. | ||
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
| // REQUIRES: matrix-xmx8 | ||
|
|
||
| // RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=1 | ||
| // RUN: %CPU_RUN_PLACEHOLDER %t.out | ||
| // RUN: %GPU_RUN_PLACEHOLDER %t.out | ||
|
|
||
| #include <iostream> | ||
| #include <random> | ||
| #include <sycl/sycl.hpp> | ||
|
|
||
| using namespace sycl; | ||
| using namespace sycl::ext::oneapi::experimental::matrix; | ||
| using bfloat16 = sycl::ext::oneapi::bfloat16; | ||
|
|
||
| #define SG_SZ 8 | ||
|
|
||
| #include "../joint_matrix_apply_bf16_impl.hpp" | ||
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,24 @@ | ||
| //==----------- joint_matrix_apply_bf16.cpp - DPC++ joint_matrix-----------==// | ||
| // | ||
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
| // See https://llvm.org/LICENSE.txt for license information. | ||
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
| // REQUIRES: matrix | ||
|
|
||
| // RUN: %clangxx -fsycl %s -o %t.out -DSYCL_EXT_ONEAPI_MATRIX_VERSION=4 | ||
| // RUN: %CPU_RUN_PLACEHOLDER %t.out | ||
| // RUN: %GPU_RUN_PLACEHOLDER %t.out | ||
|
|
||
| #include <iostream> | ||
| #include <random> | ||
| #include <sycl/sycl.hpp> | ||
|
|
||
| using namespace sycl; | ||
| using namespace sycl::ext::oneapi::experimental::matrix; | ||
| using bfloat16 = sycl::ext::oneapi::bfloat16; | ||
|
|
||
| #define SG_SZ 16 | ||
|
|
||
| #include "joint_matrix_apply_bf16_impl.hpp" |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,95 @@ | ||
|
|
||
| #define TM 8 | ||
| #define TN SG_SZ | ||
| #define TK 16 | ||
|
|
||
| static float make_fp32(bfloat16 x) { | ||
| unsigned int y = sycl::bit_cast<uint16_t>(x); | ||
| y = y << 16; | ||
| float *res = reinterpret_cast<float *>(&y); | ||
| return *res; | ||
| } | ||
|
|
||
| template <typename T, size_t NUM_ROWS, size_t NUM_COLS> struct big_matrix { | ||
| public: | ||
| T *mat; | ||
|
|
||
| public: | ||
| T *get_data() { return mat; } | ||
| void set_data(T *data) { mat = data; } | ||
| big_matrix(T *data) : mat(data) {} | ||
| }; | ||
|
|
||
| template <typename T> struct apply_add { | ||
| void operator()(T &x) const { x = x + bfloat16(2); } | ||
| }; | ||
|
|
||
| template <typename T, size_t M, size_t N, typename F> | ||
| void matrix_verify_add(queue q, big_matrix<T, M, N> &A, nd_range<2> &r, | ||
| const float ref, F &&lambda) { | ||
| buffer<bfloat16, 2> bufA(A.get_data(), range<2>(M, N)); | ||
|
|
||
| q.submit([&](handler &cgh) { | ||
| accessor accA{bufA, cgh}; | ||
|
|
||
| cgh.parallel_for(r, [accA, lambda]( | ||
| nd_item<2> spmd_item) [[intel::reqd_sub_group_size( | ||
| SG_SZ)]] { | ||
| const auto global_idx = spmd_item.get_global_id(0); | ||
| const auto global_idy = spmd_item.get_global_id(1); | ||
| const auto sg_startx = global_idx - spmd_item.get_local_id(0); | ||
| const auto sg_starty = global_idy - spmd_item.get_local_id(1); | ||
|
|
||
| sub_group sg = spmd_item.get_sub_group(); | ||
| joint_matrix<sub_group, T, use::a, TM, TK, layout::row_major> sub_a; | ||
|
|
||
| joint_matrix_fill(sg, sub_a, bfloat16(5.0)); | ||
|
|
||
| joint_matrix_apply(sg, sub_a, lambda); | ||
|
|
||
| ext::intel::experimental::matrix::joint_matrix_store( | ||
| sg, sub_a, | ||
| accA.get_pointer() + (sg_startx * TM) * N + sg_starty / SG_SZ * TN, | ||
| N); | ||
| }); // parallel for | ||
| }).wait(); | ||
| // Check if the results are correct | ||
| { | ||
| host_accessor Acc{bufA}; | ||
| assert(std::all_of(Acc.begin(), Acc.end(), [=](auto Elem) { | ||
| return (std::fabs(static_cast<float>(make_fp32(Elem) - ref)) < | ||
| std::numeric_limits<float>::epsilon()); | ||
| })); | ||
| } | ||
| } | ||
|
|
||
| static constexpr size_t MATRIX_M = TM * 2; | ||
| static constexpr size_t MATRIX_N = TN * 2; | ||
| bfloat16 A[MATRIX_M][MATRIX_N]; | ||
| float D[MATRIX_M][MATRIX_N]; | ||
|
|
||
| void matrix_ops_ref(float *D, int M, int N) { | ||
| for (int m = 0; m < M; m++) | ||
| for (int n = 0; n < N; n++) { | ||
| *(D + m * N + n) = 0; | ||
| *(D + m * N + n) *= 2; | ||
| } | ||
| } | ||
|
|
||
| int main() { | ||
|
|
||
| big_matrix<float, MATRIX_M, MATRIX_N> MD((float *)&D); | ||
| big_matrix<bfloat16, MATRIX_M, MATRIX_N> MA((bfloat16 *)&A); | ||
|
|
||
| size_t NDRangeM = MATRIX_M / TM; | ||
| size_t NDRangeN = MATRIX_N / TN; | ||
| queue q; | ||
| nd_range<2> r({NDRangeM, NDRangeN * SG_SZ}, {1, 1 * SG_SZ}); | ||
|
|
||
| matrix_verify_add<bfloat16, MATRIX_M, MATRIX_N>( | ||
| q, MA, r, 7.0, [=](bfloat16 &x) { x = x + bfloat16(2); }); | ||
| matrix_verify_add<bfloat16, MATRIX_M, MATRIX_N>(q, MA, r, 7.0, | ||
| apply_add<bfloat16>()); | ||
| std::cout << "Passed\n"; | ||
| return 0; | ||
| } |
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Uh oh!
There was an error while loading. Please reload this page.