Skip to content

[SYCL] Update Matrix tests to use access_mode #22824

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
sarnex merged 2 commits into from
Jul 30, 2026
Merged

[SYCL] Update Matrix tests to use access_mode #22824

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
d335cea
[SYCL] Update Matrix tests to use access_mode
KornevNikita Jul 30, 2026
9efde83
format
KornevNikita Jul 30, 2026
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
2 changes: 1 addition & 1 deletion sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_half_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,7 +12,7 @@ template <size_t TileRows, size_t TileCols> class divide;
template <size_t TileRows, size_t TileCols> class logic;

template <typename T, size_t Rows, size_t Cols, typename R>
void assert_ops_ref(host_accessor<T, 2, access::mode::read> C, const R ref) {
void assert_ops_ref(host_accessor<T, 2, access_mode::read> C, const R ref) {
for (size_t i = 0; i < Rows; i++)
for (size_t j = 0; j < Cols; j++) {
auto diff = C[i][j] - ref;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,7 +8,7 @@
//===----------------------------------------------------------------------===//

template <typename T, size_t NUM_ROWS, size_t NUM_COLS>
void assert_ops_ref(host_accessor<T, 2, access::mode::read> mat,
void assert_ops_ref(host_accessor<T, 2, access_mode::read> mat,
const float ref) {
for (size_t i = 0; i < NUM_ROWS; i++)
for (size_t j = 0; j < NUM_COLS; j++) {
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_int8_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@ template <size_t TileRows, size_t TileCols> class divide;
template <size_t TileRows, size_t TileCols> class logic;

template <typename T, size_t Rows, size_t Cols, typename R>
void assert_ops_ref(host_accessor<T, 2, access::mode::read> C, const R ref) {
void assert_ops_ref(host_accessor<T, 2, access_mode::read> C, const R ref) {
for (size_t i = 0; i < Rows; i++)
for (size_t j = 0; j < Cols; j++) {
auto diff = C[i][j] - ref;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_int8_packed_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@ template <size_t TileRows, size_t TileCols> class divide;
template <size_t TileRows, size_t TileCols> class logic;

template <typename T, size_t Rows, size_t Cols, typename TResult>
void assert_ops_ref(host_accessor<T, 2, access::mode::read> C,
void assert_ops_ref(host_accessor<T, 2, access_mode::read> C,
const TResult ref) {
for (size_t i = 0; i < Rows; i++)
for (size_t j = 0; j < Cols; j++) {
Expand Down
3 changes: 1 addition & 2 deletions sycl/test-e2e/Matrix/Inputs/element_wise_all_ops_tf32_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,8 +11,7 @@
#define TK 8

template <typename T, size_t M, size_t N>
void assert_ops_ref(host_accessor<T, 2, access::mode::read> C,
const float ref) {
void assert_ops_ref(host_accessor<T, 2, access_mode::read> C, const float ref) {
for (size_t i = 0; i < M; i++)
for (size_t j = 0; j < N; j++) {
auto diff = C[i][j] - ref;
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Matrix/Inputs/element_wise_all_sizes_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,7 +11,7 @@ static constexpr size_t M_MULTIPLIER = 16;
template <typename T, size_t TileM, size_t TileN, size_t TileK> class add;

template <typename T, size_t M, size_t N>
void assert_ops_ref(host_accessor<T, 2, access::mode::read_write> C,
void assert_ops_ref(host_accessor<T, 2, access_mode::read_write> C,
const T ref) {
for (size_t i = 0; i < M; i++)
for (size_t j = 0; j < N; j++) {
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/Matrix/Inputs/element_wise_ops_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,9 +21,9 @@ void matrix_multiply(big_matrix<Tc, M, N> &C, big_matrix<Ta, M, K> &A,
queue q;
size_t sg_size = get_sg_size<kernel_name>(q);
q.submit([&](handler &cgh) {
auto accC = bufC.template get_access<access::mode::read_write>(cgh);
auto accA = bufA.template get_access<access::mode::read_write>(cgh);
auto accB = bufB.template get_access<access::mode::read_write>(cgh);
auto accC = bufC.template get_access<access_mode::read_write>(cgh);
auto accA = bufA.template get_access<access_mode::read_write>(cgh);
auto accB = bufB.template get_access<access_mode::read_write>(cgh);

cgh.parallel_for<kernel_name>(
nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}),
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/Matrix/Inputs/joint_matrix_apply_cuda.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,7 +38,7 @@ void matrix_verify_lambda(queue q,
buffer<T2, 2> bufC(C.get_data(), range<2>(N * nWGperDim, M * nWGperDim));

q.submit([&](handler &cgh) {
accessor<T2, 2, access::mode::read_write, target::device> accC(bufC, cgh);
accessor<T2, 2, access_mode::read_write, target::device> accC(bufC, cgh);

cgh.parallel_for<KernelApply<T, T2, M, K, N>>(
r, [accC, lambda](
Expand Down
4 changes: 2 additions & 2 deletions sycl/test-e2e/Matrix/Inputs/joint_matrix_down_convert_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,8 +26,8 @@ void matrix_copy(big_matrix<T1, M, N> &C, big_matrix<T2, M, K> &A) {
queue q;
size_t sg_size = get_sg_size<class copy>(q);
q.submit([&](handler &cgh) {
auto accC = bufC.get_access<access::mode::read_write>(cgh);
auto accA = bufA.get_access<access::mode::write>(cgh);
auto accC = bufC.get_access<access_mode::read_write>(cgh);
auto accA = bufA.get_access<access_mode::write>(cgh);

cgh.parallel_for<class copy>(
nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}),
Expand Down
16 changes: 8 additions & 8 deletions sycl/test-e2e/Matrix/Inputs/joint_matrix_gemm_cuda.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -113,8 +113,8 @@ void test(queue &q) {
buffer<bfloat16, 1> bufA(A, range<1>(Big_M * Big_K));
buffer<bfloat16, 1> bufB(B, range<1>(Big_K * Big_N));
q.submit([&](handler &cgh) {
accessor<bfloat16, 1, access::mode::write, target::device> accA(bufA,
cgh);
accessor<bfloat16, 1, access_mode::write, target::device> accA(bufA,
cgh);

cgh.parallel_for<KernelName<Tm, Tc, class copyA, M, K, N, layout_A,
layout_B, layout_C>>(
Expand All @@ -124,8 +124,8 @@ void test(queue &q) {
});
});
q.submit([&](handler &cgh) {
accessor<bfloat16, 1, access::mode::write, target::device> accB(bufB,
cgh);
accessor<bfloat16, 1, access_mode::write, target::device> accB(bufB,
cgh);

cgh.parallel_for<KernelName<Tm, Tc, class copyB, M, K, N, layout_A,
layout_B, layout_C>>(
Expand All @@ -142,10 +142,10 @@ void test(queue &q) {
buffer<Td, 1> bufD(D, range<1>(Big_M * Big_N));

q.submit([&](handler &cgh) {
accessor<Tm, 1, access::mode::read, target::device> accA(bufA, cgh);
accessor<Tm, 1, access::mode::read, target::device> accB(bufB, cgh);
accessor<Tc, 1, access::mode::read, target::device> accC(bufC, cgh);
accessor<Td, 1, access::mode::write, target::device> accD(bufD, cgh);
accessor<Tm, 1, access_mode::read, target::device> accA(bufA, cgh);
accessor<Tm, 1, access_mode::read, target::device> accB(bufB, cgh);
accessor<Tc, 1, access_mode::read, target::device> accC(bufC, cgh);
accessor<Td, 1, access_mode::write, target::device> accD(bufD, cgh);

range<2> LocalRange = {1, N_THREADS_PER_MATRIX_OP};
range<2> GlobalRange = {Sub_Tiles_M,
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/Matrix/Inputs/joint_matrix_int8_colmajorA_colmajorB_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,9 +28,9 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
queue q;
size_t sg_size = get_sg_size<class imatrix>(q);
q.submit([&](handler &cgh) {
auto accC = bufC.get_access<access::mode::read_write>(cgh);
auto accA = bufA.get_access<access::mode::read_write>(cgh);
auto accB = bufB.get_access<access::mode::read_write>(cgh);
auto accC = bufC.get_access<access_mode::read_write>(cgh);
auto accA = bufA.get_access<access_mode::read_write>(cgh);
auto accB = bufB.get_access<access_mode::read_write>(cgh);

cgh.parallel_for<class imatrix>(
nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}),
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/Matrix/Inputs/joint_matrix_tf32_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,9 +30,9 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,
queue q;
size_t sg_size = get_sg_size<class imatrix>(q);
q.submit([&](handler &cgh) {
auto accC = bufC.get_access<access::mode::read_write>(cgh);
auto accA = bufA.get_access<access::mode::read_write>(cgh);
auto accB = bufB.get_access<access::mode::read_write>(cgh);
auto accC = bufC.get_access<access_mode::read_write>(cgh);
auto accA = bufA.get_access<access_mode::read_write>(cgh);
auto accB = bufB.get_access<access_mode::read_write>(cgh);

cgh.parallel_for<class imatrix>(
nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}),
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/Matrix/elemwise_irreg_size_ops_bf16.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,9 +47,9 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,

queue q;
q.submit([&](handler &cgh) {
auto accC = bufC.get_access<access::mode::read_write>(cgh);
auto accA = bufA.get_access<access::mode::read_write>(cgh);
auto accB = bufB.get_access<access::mode::read_write>(cgh);
auto accC = bufC.get_access<access_mode::read_write>(cgh);
auto accA = bufA.get_access<access_mode::read_write>(cgh);
auto accB = bufB.get_access<access_mode::read_write>(cgh);

cgh.parallel_for<class imatrix>(
nd_range<2>({NDRangeM, NDRangeN * SG_SZ}, {1, 1 * SG_SZ}),
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/Matrix/joint_matrix_16bit_colmajorA_colmajorB.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,9 +41,9 @@ void matrix_multiply(big_matrix<T1, M, N> &C, big_matrix<T2, M, K> &A,
std::cout << "subgroup size " << sg_size << " ";

q.submit([&](handler &cgh) {
auto accC = bufC.get_access<access::mode::read_write>(cgh);
auto accA = bufA.template get_access<access::mode::read_write>(cgh);
auto accB = bufB.template get_access<access::mode::read_write>(cgh);
auto accC = bufC.get_access<access_mode::read_write>(cgh);
auto accA = bufA.template get_access<access_mode::read_write>(cgh);
auto accB = bufB.template get_access<access_mode::read_write>(cgh);

cgh.parallel_for<class imatrix<T2, TM, TN, TK>>(
nd_range<2>({NDRangeM, NDRangeN * sg_size}, {1, 1 * sg_size}),
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/Matrix/joint_matrix_query_default.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,9 +62,9 @@ void matrix_multiply(big_matrix<T1, NUM_ROWS_C, NUM_COLS_C> &C,

queue q;
q.submit([&](handler &cgh) {
auto accC = bufC.get_access<access::mode::read_write>(cgh);
auto accA = bufA.get_access<access::mode::read_write>(cgh);
auto accB = bufB.get_access<access::mode::read_write>(cgh);
auto accC = bufC.get_access<access_mode::read_write>(cgh);
auto accA = bufA.get_access<access_mode::read_write>(cgh);
auto accB = bufB.get_access<access_mode::read_write>(cgh);

cgh.parallel_for<class imatrix>(
nd_range<2>({NDRangeM, NDRangeN * SG_SZ}, {1, 1 * SG_SZ}),
Expand Down
Loading