Intel gpu backend gemm pipeline #89

Jiaxingla · 2024-07-03T08:09:59Z

High performance gemm pipeline for pvc.
Enable the prefetch by copy atom.
Try the epilogue like ReLU and Softmax.

examples/sycl/pvc/pvc_bfloat_dpas_gemm_cute.cpp

tools/util/include/cutlass/util/reference/device/tensor_compare.h

include/cutlass/epilogue/collective/default_epilogue.hpp

rolandschulz · 2024-07-10T17:21:25Z

include/cute/arch/copy_xe.hpp

@@ -44,6 +44,18 @@ namespace cute
  inline x { assert(false); }
 #endif

+enum LSC_LDCC {


https://github.com/NVIDIA/cutlass/blob/main/media/docs/programming_guidelines.md#use-scoped-enums

examples/sycl/pvc/pvc_bfloat_dpas_gemm_cute.cpp

rolandschulz · 2024-07-10T17:39:08Z

examples/sycl/pvc/pvc_bfloat_dpas_gemm_cute.cpp

  //
  // Methods
  //

-  bool verify(const ProblemShapeType& problem_size, ElementCompute alpha, ElementCompute beta) {
+  bool verify(const ProblemShapeType &problem_size, ElementCompute alpha,


Did you use automatic code formatting? https://github.com/NVIDIA/cutlass/blob/main/media/docs/programming_guidelines.md#no-automatic-code-formatting

I hope we can use it because I think having to review formatting is a pain. But if so we need to find a tool which follows all of https://github.com/NVIDIA/cutlass/blob/main/media/docs/programming_guidelines.md and only apply to lines we change (e.g. git clang-format) so that we don't reformat any of the upstream code.

Here format is again wrong (west const, space alignment, line length). Please try to fix it everywhere and I'll review formatting after.

--Did you use automatic code formatting?
Yes, we used a wrong code formatting.
I'll correct the clang-format and try my best to fix the format everywhere with cutlass guidelines.

if you are using clang-format please add the config

we aren't using Chang-format?

examples/sycl/pvc/pvc_bfloat_dpas_gemm_cute.cpp

rolandschulz · 2024-07-10T18:09:00Z

include/cute/arch/copy_xe.hpp

+        template <class T>
+        CUTE_HOST_DEVICE static void copy(const void *baseoffset, int width,
+                int height, int pitch, intel::coord_t coord) {
+#if defined(SYCL_INTEL_TARGET)


This pattern is the same as on the unmodified line 170.

I think we can improve this by changing __SYCL_DEVICE_ONLY__ to

#ifdef __SYCL_DEVICE_ONLY__ #ifdef SYCL_INTEL_TARGET #define SYCL_DEVICE_BUILTIN(x) SYCL_EXTERNAL extern "C" x #else #define SYCL_DEVICE_BUILTIN(x) \ inline x { CUTE_INVALID_CONTROL_PATH("Trying to use IGC built-in on non-Intel hardware"); } #endif #else #define SYCL_DEVICE_BUILTIN(x) \ inline x { CUTE_INVALID_CONTROL_PATH("Trying to use device built-in on host."); } #endif

Up to you whether you want to do this as part of this PR.

include/cutlass/gemm/collective/intel_pvc_mma.hpp

include/cutlass/relatively_equal.h

include/cutlass/gemm/collective/intel_pvc_mma.hpp

include/cutlass/gemm/kernel/intel_pvc_gemm.hpp

examples/sycl/pvc/pvc_gemm.cpp

mehdi-goli · 2024-07-30T08:25:42Z

examples/sycl/pvc/pvc_gemm.cpp

@@ -362,14 +340,14 @@ int main(int argc, const char** argv)
  using LayoutD = cutlass::layout::RowMajor;

  using GmemTiledCopyA = XE_2D_U16x8x16x4x2_LD_N;


adding some explanation for naming conventions of copy function

mehdi-goli · 2024-07-30T09:17:42Z

include/cute/atom/copy_atom.hpp

+    } else if constexpr (is_tuple<typename Tensor<SEngine,
+                                 SLayout>::engine_type::iterator::
+                                         value_type>::value) {
+        return copy_unpack(*this, src, dst);


As discussed please remove this elseif condition and make the changes in the copy traits definition so the execution go into the first if check

For this you need to make SrcLayout in the prefetch copy trait to match decltype(size(src))>::value

include/cute/arch/mma_xe.hpp

mehdi-goli · 2024-07-30T09:22:39Z

examples/sycl/pvc/pvc_gemm_with_epilogue_relu.cpp

Please remove the VNNI file from here as it is not needed https://github.com/codeplaysoftware/cutlass-fork/pull/89/files#diff-eef47671b10b1dafabc54befe9bfbc0ed6a690a6688bf0fc638c7f2bf2b72ebcL61-R75

Co-authored-by: Mehdi Goli <[email protected]>

muhammad-tanvir-1211 · 2024-07-31T12:25:05Z

include/cutlass/gemm/collective/intel_pvc_mma.hpp

-   }
+    int prefetch_k = 0;
+
+    // Manually set the value to 1


Shouldn't this comment say Manually set the value to 3?

aacostadiaz

Thanks for addressing the comments. LGTM!

mehdi-goli

LGTM!

rolandschulz · 2024-07-31T21:32:50Z

Some of my comments haven't been addressed. Are you planning to address them before merging or after?

Jiaxingla · 2024-08-01T01:53:12Z

Some of my comments haven't been addressed. Are you planning to address them before merging or after?

Hi @rolandschulz , i do want to address or answer all your comments before merging. Because PR is completely different with beginning, if i have missed any of your comments, please let me know.
Apologize for any inconvenience caused.

rolandschulz · 2024-08-01T05:40:38Z

I clicked "resolve" on everything which I can see being addressed. Please check the remaining 3.

Jiaxingla · 2024-08-01T06:14:39Z

I clicked "resolve" on everything which I can see being addressed. Please check the remaining 3.

I resolved the format of enum class and answered the questions about clang-format.
The last about SYCL macro will fix in the next PR about Cute feature.

rolandschulz · 2024-08-01T07:30:41Z

include/cute/arch/copy_xe.hpp

@@ -44,6 +44,17 @@ namespace cute
  inline x { assert(false); }
 #endif

+enum CacheControl {


shouldn't this be a scoped enum?

include/cute/arch/copy_xe.hpp

rolandschulz · 2024-08-01T07:40:59Z

include/cute/arch/copy_xe.hpp

+                int height, int pitch, intel::coord_t coord) {
+#if defined(SYCL_INTEL_TARGET)
+            static_assert(sizeof(T) == 2, "Expected T to have size 2");
+            __builtin_IB_subgroup_block_read_prefetch_u16_m16k16v1(


why does this not match the non-prefetch in the same struct?

Seems to still not match (u16/u32).

aacostadiaz · 2024-08-01T16:30:20Z

CI is failing due to changes introduced in the latest nightly version of DPCPP. Rebasing on sycl_develop should resolve this.

rolandschulz · 2024-08-01T17:01:31Z

include/cute/arch/copy_xe.hpp

 };

+/// @brief This function loads data from 2D memory surface.
+/// Loads an array of rectangular regions coord(X,Y)..coord(X+W,Y+H) from global memory into registers.
+/// Loads 1x1 memory blocks, and each block size is 8x16x32bits


why do we use different ordering of dimensions for the comment and the name of the struct?

aacostadiaz · 2024-08-02T08:59:01Z

I think all comments have been addressed. If there is something missing, please let us know and it will be added in the follow up PR.

Enable the prefetch by copy atom. --------- Co-authored-by: Mehdi Goli <[email protected]>

Jiaxingla requested review from taozha2 and jiyang1011 July 3, 2024 08:10

Jiaxingla marked this pull request as ready for review July 4, 2024 05:55

Jiaxingla requested review from mehdi-goli and rolandschulz July 4, 2024 05:55

rolandschulz reviewed Jul 9, 2024

View reviewed changes

Jiaxingla force-pushed the intel_gpu_backend_pipeline branch from 189753c to d82b5a4 Compare July 10, 2024 01:28

rolandschulz reviewed Jul 10, 2024

View reviewed changes

rolandschulz requested review from muhammad-tanvir-1211 and aacostadiaz July 10, 2024 19:08

rolandschulz reviewed Jul 11, 2024

View reviewed changes

Jiaxingla force-pushed the intel_gpu_backend_pipeline branch 5 times, most recently from e7381bc to 26d8a88 Compare July 16, 2024 08:15

aacostadiaz force-pushed the sycl-develop branch from 039211b to 8f3bd67 Compare July 16, 2024 10:06

Jiaxingla added 11 commits July 16, 2024 23:01

apply patch of gemm pipeline

07f36e5

fix format of copyright

d4cf3eb

replace the macro of cache flush and idx

665f9be

auto format

59c0ce4

auto format

bdadf1e

fix comments about prefetch

9e23cd6

fix comments of enum and sycl macro

60adb24

update from tensor library repo

1e3f855

fix format

8e951d1

rm redundancy code

c92adb3

resolve conflict

6bdda75

Jiaxingla force-pushed the intel_gpu_backend_pipeline branch from 576261d to 6bdda75 Compare July 17, 2024 06:31

revert the change of nv hpp

3496593

modify comments

c97ccd8

mehdi-goli reviewed Jul 30, 2024

View reviewed changes

examples/sycl/pvc/pvc_gemm.cpp Outdated Show resolved Hide resolved

mehdi-goli reviewed Jul 30, 2024

View reviewed changes

include/cute/arch/mma_xe.hpp Outdated Show resolved Hide resolved

mehdi-goli reviewed Jul 30, 2024

View reviewed changes

Jiaxingla and others added 4 commits July 31, 2024 14:10

Update pvc_gemm

ede5c03

Co-authored-by: Mehdi Goli <[email protected]>

Update mma_xe

f9aae6f

Co-authored-by: Mehdi Goli <[email protected]>

more comments

7878a7c

code clean

4c42645

muhammad-tanvir-1211 reviewed Jul 31, 2024

View reviewed changes

Jiaxingla and others added 2 commits July 31, 2024 20:34

fix typo

abbbe4f

revert the change of copy_atom

8e9a84f

aacostadiaz approved these changes Jul 31, 2024

View reviewed changes

mehdi-goli approved these changes Jul 31, 2024

View reviewed changes

rename enum of LSC_LDCC

ea30c83

rolandschulz reviewed Aug 1, 2024

View reviewed changes

Jiaxingla added 2 commits August 1, 2024 01:00

fix typo

043fbea

scope enums

abf38bd

rolandschulz reviewed Aug 1, 2024

View reviewed changes

Jiaxingla added 2 commits August 1, 2024 18:07

modify commment of copy

5193329

remove useless copy

b854995

aacostadiaz merged commit 0b5c911 into codeplaysoftware:sycl-develop Aug 2, 2024
4 checks passed

aacostadiaz pushed a commit to aacostadiaz/cutlass-fork that referenced this pull request Aug 6, 2024

Intel gpu backend gemm pipeline (codeplaysoftware#89)

a1a943f

Enable the prefetch by copy atom. --------- Co-authored-by: Mehdi Goli <[email protected]>

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Intel gpu backend gemm pipeline #89

Intel gpu backend gemm pipeline #89

Jiaxingla commented Jul 3, 2024 •

edited

Loading

rolandschulz Jul 10, 2024

rolandschulz Jul 10, 2024

Jiaxingla Jul 11, 2024

rolandschulz Jul 15, 2024

rolandschulz Aug 1, 2024

rolandschulz Jul 10, 2024

mehdi-goli Jul 30, 2024

mehdi-goli Jul 30, 2024

aacostadiaz Jul 31, 2024

mehdi-goli Jul 30, 2024

muhammad-tanvir-1211 Jul 31, 2024

aacostadiaz left a comment

mehdi-goli left a comment

rolandschulz commented Jul 31, 2024

Jiaxingla commented Aug 1, 2024 •

edited

Loading

rolandschulz commented Aug 1, 2024

Jiaxingla commented Aug 1, 2024

rolandschulz Aug 1, 2024

rolandschulz Aug 1, 2024

rolandschulz Aug 1, 2024

aacostadiaz commented Aug 1, 2024

rolandschulz Aug 1, 2024

aacostadiaz commented Aug 2, 2024

		@@ -362,14 +340,14 @@ int main(int argc, const char** argv)
		using LayoutD = cutlass::layout::RowMajor;

		using GmemTiledCopyA = XE_2D_U16x8x16x4x2_LD_N;

Intel gpu backend gemm pipeline #89

Intel gpu backend gemm pipeline #89

Conversation

Jiaxingla commented Jul 3, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

aacostadiaz left a comment

Choose a reason for hiding this comment

mehdi-goli left a comment

Choose a reason for hiding this comment

rolandschulz commented Jul 31, 2024

Jiaxingla commented Aug 1, 2024 • edited Loading

rolandschulz commented Aug 1, 2024

Jiaxingla commented Aug 1, 2024

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

aacostadiaz commented Aug 1, 2024

Choose a reason for hiding this comment

aacostadiaz commented Aug 2, 2024

Jiaxingla commented Jul 3, 2024 •

edited

Loading

Jiaxingla commented Aug 1, 2024 •

edited

Loading