Skip to content

[SYCL][Matrix] Add spec document for the matrix extension interface and its first implementation for AMX #3551

New issue

Have a question about this project? # for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “#”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? # to your account

Merged
merged 11 commits into from
Jul 2, 2021

Conversation

dkhaldi
Copy link
Contributor

@dkhaldi dkhaldi commented Apr 15, 2021

Signed-off-by: Dounia dounia.khaldi@intel.com

@dkhaldi dkhaldi requested a review from a team as a code owner April 15, 2021 16:53
@bader bader added the spec extension All issues/PRs related to extensions specifications label Apr 15, 2021
@dkhaldi dkhaldi requested a review from Pennycook April 15, 2021 16:57
@keryell
Copy link
Contributor

keryell commented Apr 15, 2021

That seems quite interesting.
It is not clear to me in which namespace they are defined.
Probably they could be defined as friend in the matrix class itself so the operations can be found by ADL, avoiding to type the namespace again and again everywhere...
It would be interesting to investigate whether providing also member functions would simplify the API.
Provide both so it is possible to use the best one for each use case, while waiting for https://en.wikipedia.org/wiki/Uniform_Function_Call_Syntax to land into C++?

@dkhaldi
Copy link
Contributor Author

dkhaldi commented Apr 16, 2021

That seems quite interesting.
It is not clear to me in which namespace they are defined.

@keryell They are defined cl::sycl::ext::intel::matrix;
I will add the "using namespace cl::sycl::ext::intel::matrix;"
to the code example

Probably they could be defined as friend in the matrix class itself so the operations can be found by ADL, avoiding to retype the namespace everywhere...
It would be interesting to investigate whether providing also member functions would simplify the API.

They are free functions to match the group library algorithms. But I will add your comment to the open questions section.

…d its first implementation for AMX

Signed-off-by: Dounia <dounia.khaldi@intel.com>
@keryell
Copy link
Contributor

keryell commented Apr 17, 2021

I will add the "using namespace cl::sycl::ext::intel::matrix;"

On the long term, using either member functions or ADL could avoid this requirement that adds to the namespace clutter...

…espace sycl::ext::intel::experimental::matrix, remove the C++ specific API from the document, better formatting

Signed-off-by: Dounia <dounia.khaldi@intel.com>
@dkhaldi
Copy link
Contributor Author

dkhaldi commented Jun 8, 2021

ping @gmlueck

…escription to a different section towards the end

Signed-off-by: Dounia <dounia.khaldi@intel.com>
dkhaldi added 3 commits June 14, 2021 13:53
…espace, some formatting and rewording

Signed-off-by: Dounia <dounia.khaldi@intel.com>
- Add the implementation status to doc/extensions/README.md
- Add "API description vs what is actually implemented" question
like dynamic_ extent and Group to the open questions
- Add more clarification about packed_a and packed_b layout,
and difference between layouts on matrix and in load/store functions

Signed-off-by: Dounia <dounia.khaldi@intel.com>
…tation capabilities in each of the API sections, and some improvement to the layouts explanation text
Copy link
Contributor

@gmlueck gmlueck left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is looking a lot better! Now that I understand the intent, I have some more comments.

dkhaldi added 2 commits July 1, 2021 13:08
- Remove all the comments from the asciidoc file 
- Add more explicitly in two other places that "a kernel using this extension must be decorated with the [[sycl::reqd_sub_group_size(1)]] attribute"
- Add 8 bit example for VNNI transform
- Incorporate rewording suggestions from Greg.
@bader bader merged commit ace4c73 into intel:sycl Jul 2, 2021
alexbatashev added a commit to alexbatashev/llvm that referenced this pull request Jul 2, 2021
* upstream/sycl: (649 commits)
  [SYCL][Driver][NFC] Update integration footer test for 32-bit host (intel#4039)
  [SYCL][L0] Initialize descriptor .stype and .pNext (intel#4032)
  [SYCL] Add sycl::kernel::get_kernel_bundle method (intel#3855)
  [SYCL] Add support for device UUID as a SYCL extension. (intel#3696)
  [SYCL][Matrix] Add spec document for the matrix extension interface and its first implementation for AMX (intel#3551)
  Fix debug build mangler test after PR#3992 (8f38045). (intel#4033)
  [Driver][SYCL] Restrict user -include file in final integration footer step (intel#4036)
  [SYCL] [Tests] Do not copy device binary image mocks (intel#4023)
  [SYCL][Doc] Update docs to reflect new compiler features (intel#4030)
  [SYCL][CUDA] cl_khr_fp16 extension connected to cuda PI. (intel#4029)
  [SYCL][NFC] Refactor RT unit tests (intel#4021)
  [SYCL] Switch to using integration footer by default (intel#3777)
  [SYCL][CUDA] Add the Use Default Stream property (intel#4004)
  Uplift GPU RT version for Linux to 21.24.20098 (intel#4003)
  [SYCL][CUDA] atomic_ref.fetch_add used for fp64 reduction if device.has(atomic64) (intel#3950)
  [Driver][SYCL] Differentiate host dependency link from regular host link (intel#4002)
  [SYCL][ESIMD] Support device half type in intrinsics. (intel#4024)
  [SYCL] Allow fpga_reg only for PODs and Trivially-copyable structs (intel#3643)
  [SYCL][FPGA] Restore legacy debug info version for the hardware (intel#3991)
  [SYCL][PI][L0] Force reset of memcpy command-list. (intel#4001)
  ...
# for free to join this conversation on GitHub. Already have an account? # to comment
Labels
spec extension All issues/PRs related to extensions specifications
Projects
None yet
Development

Successfully merging this pull request may close these issues.

6 participants