[SYCL][matrix] add a new test for irregular slicing on packed matrix … #899
Conversation
…that calculates sum of rows of the matrix
|
@Nuullll, this is a test for irregular slicing but it is still a draft. |
Thanks! |
| const auto sg_starty = global_idy - spmd_item.get_local_id(1); | ||
|
|
||
| ext::oneapi::sub_group sg = spmd_item.get_sub_group(); | ||
| joint_matrix<T, TK, TN, matrix_layout::packed_b> sub_b(sg); |
There was a problem hiding this comment.
IIUC, the template arguments specify the logical size of sub_b, then why is it a 32 x 8 matrix? Shouldn't it be 8 x 8?
There was a problem hiding this comment.
I will correct the test, it should be 32x8 as I am imitating a real use for matrix B
| for (int i = 0; i < MATRIX_M / 4; i++) { | ||
| for (int j = 0; j < MATRIX_N * 4; j++) { | ||
| B[i][j] = i; |
There was a problem hiding this comment.
The loop size doesn't match B's declaration size at line 99. We'll end up with:
// 16x16 view
B = [0, 0, 0, ... 0;
1, 1, 1, ... 1;
2, 2, 2, ... 2;
3, 3, 3, ... 3;
3, 3, 3, ... 3;
3, 3, 3, ... 3;
3, 3, 3, ... 3;
0, 0, 0, ... 0;
...
0, 0, 0, ... 0;]
Is this expected?
| // TODO: communicate this mapping information to the user using a | ||
| // query interface |
There was a problem hiding this comment.
Should this TODO be here? Is it expected to affect the test when it is implemented?
There was a problem hiding this comment.
Only half of the TODO was removed.
| void matrix_sum_rows(queue q, big_matrix<T, M, N> &B, nd_range<2> &r) { | ||
| buffer<int8_t, 2> bufB(B.get_data(), range<2>(M, N)); | ||
| // size of vector is known because SG size of set by the user in this case | ||
| buffer<int> sum_rows_v(M); // there are total of tK/4 * 2, 16 rows |
There was a problem hiding this comment.
Initialize the buffer with zeros, otherwise atomic_fetch_add would add numbers to uninitialized values and get weird results.
| buffer<int> sum_rows_v(M); // there are total of tK/4 * 2, 16 rows | |
| int sum_rows[M] = {0}; | |
| buffer<int> sum_rows_v(sum_rows, M); // there are total of tK/4 * 2, 16 rows |
|
@steffenlarsen is SYCL :: SubGroup/attributes.cpp expected to fail? |
There are some known sporadic failures caused by driver issues. We hope to have them fixed soon. |
|
@steffenlarsen can we merge this PR? |
CMPLRTOOLS-26492: enable ESIMD_EMULATOR for optset opt_use_esimd_emu
…that calculates sum of rows of the matrix (intel#899)
…that calculates sum of rows of the matrix (intel/llvm-test-suite#899)
…that calculates sum of rows of the matrix