Skip to content

Commit

Permalink
Fix copy operation and mma tile definition (#147)
Browse files Browse the repository at this point in the history
  • Loading branch information
aacostadiaz authored Oct 25, 2024
1 parent 0162a1e commit 1dc136d
Showing 1 changed file with 7 additions and 7 deletions.
14 changes: 7 additions & 7 deletions examples/sycl/pvc/pvc_gemm_with_epilogue_relu.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -316,15 +316,15 @@ int main(int argc, const char** argv)
using LayoutC = cutlass::layout::RowMajor;
using LayoutD = cutlass::layout::RowMajor;

using GmemTiledCopyA = XE_2D_U16x8x16x4x2_LD_N;
using GmemTiledCopyB = XE_2D_U16x16x16x2x2_V;
using GmemTiledCopyA = XE_2D_U16x8x16_LD_N;
using GmemTiledCopyB = XE_2D_U16x16x16_LD_V;

// Workgroup-level tile
using TileShape = Shape<_256, _256, _32>;
using TileShape = Shape<_256, _128, _16>;

using TiledMma = TiledMMA<MMA_Atom<XE_8x16x16_F32BF16BF16F32_TT>,
Layout<Shape<_1,_1,_1>>,
Tile<_32,_64,_32>>; // Subgroup level-tile
Layout<Shape<_8,_2,_1>>,
Tile<_64,_32,_16>>; // Subgroup level-tile

constexpr int PipelineStages = 3;
using GEMMDispatchPolicy = cutlass::gemm::MainloopIntelPVC<PipelineStages>;
Expand All @@ -343,9 +343,9 @@ int main(int argc, const char** argv)
ElementOutput,
cutlass::gemm::TagToStrideC_t<LayoutD>,
FusionCallBacks,
XE_2D_U32x8x16x1x1_LD_N,
XE_2D_U32x8x16_LD_N,
void, void,
XE_2D_U32x8x16x1x1_ST_N,
XE_2D_U32x8x16_ST_N,
void, void>;

// Mainloop
Expand Down

0 comments on commit 1dc136d

Please sign in to comment.