Skip to content

Reduce memory utilization #459

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

Open
wants to merge 5 commits into
base: sycl-develop
Choose a base branch
from
+1,763 −1,755
Open

Reduce memory utilization #459

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
1deb358
Reduce memory utilization
aacostadiaz Jul 7, 2025
5d8ee86
build benchmarks
aacostadiaz Jul 7, 2025
f253665
remove space
aacostadiaz Jul 8, 2025
5b028f2
reduce threads
aacostadiaz Jul 8, 2025
1c475b6
disable benchmark
aacostadiaz Jul 8, 2025
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
9 changes: 3 additions & 6 deletions .github/workflows/intel_test.yml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -95,8 +95,9 @@ jobs:
cmake -G Ninja \
-DCUTLASS_ENABLE_SYCL=ON \
-DDPCPP_SYCL_TARGET=${{ matrix.sycl_target }} \
-DCUTLASS_SYCL_RUNNING_CI=ON
cmake --build .
-DCUTLASS_ENABLE_BENCHMARKS=OFF \
-DCUTLASS_SYCL_RUNNING_CI=ON
cmake --build . -j$(($(nproc) / 2))
- name: Unit test
shell: bash
run: |
Expand All @@ -105,7 +106,3 @@ jobs:
shell: bash
run: |
cmake --build . --target test_examples -j 1
- name: Benchmarks
shell: bash
run: |
cmake --build . --target cutlass_benchmarks
2 changes: 0 additions & 2 deletions examples/cute/tutorial/sgemm_1_sycl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -202,8 +202,6 @@ gemm_device(ProblemShape shape_MNK, CtaTiler cta_tiler,
// tAsA(i) = tAgAk(i);
// }

cp_async_fence(); // Label the end of (potential) cp.async instructions
cp_async_wait<0>(); // Sync on all (potential) cp.async instructions
syclcompat::wg_barrier();// Wait for all threads to write to smem

// Compute gemm on tC thread-partitioned smem
Expand Down
9 changes: 4 additions & 5 deletions include/cute/atom/copy_atom.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -729,11 +729,15 @@ print_latex_copy(LayoutS const& S, ThrIDS const& TS, // (m,n) -> (tid,vid) and

////////////////////////////////////////////////////////////////////////////////////////////////////

#if defined(SYCL_INTEL_TARGET)
#include <cute/atom/copy_traits_xe.hpp>
#else
#include <cute/atom/copy_traits_sm50.hpp>
#include <cute/atom/copy_traits_sm75.hpp>
#include <cute/atom/copy_traits_sm80.hpp>
#include <cute/atom/copy_traits_sm90.hpp>
#include <cute/atom/copy_traits_sm100.hpp>
#endif


// Config
Expand Down Expand Up @@ -761,9 +765,4 @@ print_latex_copy(LayoutS const& S, ThrIDS const& TS, // (m,n) -> (tid,vid) and
#include <cute/atom/copy_traits_sm100_tma.hpp>
#endif


#if defined(SYCL_INTEL_TARGET)
#include <cute/atom/copy_traits_xe.hpp>
#endif

////////////////////////////////////////////////////////////////////////////////////////////////////
10 changes: 3 additions & 7 deletions include/cute/atom/mma_atom.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,10 +36,6 @@
#include <cute/tensor_impl.hpp>
#include <cute/util/type_traits.hpp>

#if defined(CUTLASS_ENABLE_SYCL)
#include <cute/atom/mma_traits_xe.hpp>
#endif

namespace cute {

template <class... Args>
Expand Down Expand Up @@ -1184,6 +1180,9 @@ print_svg(TiledMMA<Args...> const &mma) {

////////////////////////////////////////////////////////////////////////////////////////////////////

#if defined(SYCL_INTEL_TARGET)
#include <cute/atom/mma_traits_xe.hpp>
#else
#include <cute/atom/mma_traits_sm61.hpp>
#include <cute/atom/mma_traits_sm70.hpp>
#include <cute/atom/mma_traits_sm75.hpp>
Expand All @@ -1194,8 +1193,5 @@ print_svg(TiledMMA<Args...> const &mma) {
#include <cute/atom/mma_traits_sm100.hpp>
#include <cute/atom/mma_traits_sm120.hpp>
#include <cute/atom/mma_traits_sm120_sparse.hpp>

#if defined(CUTLASS_ENABLE_SYCL)
#include <cute/atom/mma_traits_xe.hpp>
#endif
////////////////////////////////////////////////////////////////////////////////////////////////////
4 changes: 3 additions & 1 deletion include/cutlass/conv/collective/collective_builder.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -89,6 +89,8 @@ struct CollectiveBuilder {

/////////////////////////////////////////////////////////////////////////////////////////////////

#if !defined(SYCL_INTEL_TARGET)
#include "builders/sm90_gmma_builder.inl"
#include "builders/sm100_umma_builder.inl"
#include "builders/sm100_umma_builder.inl"
#endif
/////////////////////////////////////////////////////////////////////////////////////////////////
6 changes: 3 additions & 3 deletions include/cutlass/epilogue/collective/collective_builder.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -118,12 +118,12 @@ struct CallbacksBuilder<

/////////////////////////////////////////////////////////////////////////////////////////////////

#if defined(SYCL_INTEL_TARGET)
#include "builders/xe_builder.inl"
#else
#include "builders/sm90_builder.inl"
#include "builders/sm100_builder.inl"
#include "builders/sm120_builder.inl"

#if defined(SYCL_INTEL_TARGET)
#include "builders/xe_builder.inl"
#endif

#if defined(CUTLASS_ENABLE_SYCL)
Expand Down
7 changes: 4 additions & 3 deletions include/cutlass/epilogue/collective/collective_epilogue.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -57,6 +57,10 @@ class CollectiveEpilogue {
//
// Gemm
//
#if defined (SYCL_INTEL_TARGET)
#include "xe_epilogue.hpp"
#include "xe_array_epilogue.hpp"
#else
#include "default_epilogue.hpp"
#include "default_epilogue_array.hpp"
#include "epilogue_tensor_broadcast.hpp"
Expand All @@ -69,9 +73,6 @@ class CollectiveEpilogue {
#include "sm100_epilogue_array_nosmem.hpp"
#include "sm100_epilogue_tma_warpspecialized.hpp"
#include "sm100_epilogue_array_tma_warpspecialized.hpp"
#if defined (SYCL_INTEL_TARGET)
#include "xe_epilogue.hpp"
#include "xe_array_epilogue.hpp"
#endif
//
// Conv
Expand Down
6 changes: 3 additions & 3 deletions include/cutlass/gemm/collective/collective_builder.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,6 +37,9 @@
/////////////////////////////////////////////////////////////////////////////////////////////////

#include "cutlass/gemm/collective/collective_builder_decl.hpp"
#if defined(SYCL_INTEL_TARGET)
#include "cutlass/gemm/collective/builders/xe_mma_builder.inl"
#else
#include "cutlass/gemm/collective/builders/sm90_gmma_builder.inl"
#include "cutlass/gemm/collective/builders/sm90_sparse_gmma_builder.inl"
#if !defined(__CUDACC_RTC__)
Expand All @@ -51,9 +54,6 @@
#include "cutlass/gemm/collective/builders/sm120_sparse_mma_builder.inl"
#include "cutlass/gemm/collective/builders/sm120_blockscaled_sparse_mma_builder.inl"
#endif

#if defined(SYCL_INTEL_TARGET)
#include "cutlass/gemm/collective/builders/xe_mma_builder.inl"
#endif

#if defined(CUTLASS_ENABLE_SYCL)
Expand Down
17 changes: 8 additions & 9 deletions include/cutlass/gemm/collective/collective_mma.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -34,7 +34,14 @@


/////////////////////////////////////////////////////////////////////////////////////////////////

#if defined(SYCL_INTEL_TARGET)
#include "cutlass/gemm/collective/xe_mma.hpp"
#include "cutlass/gemm/collective/xe_array_mma.hpp"
#include "cutlass/gemm/collective/xe_array_mma_fp8.hpp"
#include "cutlass/gemm/collective/xe_mma_mixed_input.hpp"
#include "cutlass/gemm/collective/xe_mma_w8a8.hpp"
#include "cutlass/gemm/collective/xe_mma_fp8_scaling.hpp"
#else
#include "cutlass/gemm/collective/sm70_mma_twostage.hpp"
#include "cutlass/gemm/collective/sm80_mma_multistage.hpp"
#include "cutlass/gemm/collective/sm90_mma_multistage_gmma_ss_warpspecialized.hpp"
Expand Down Expand Up @@ -68,14 +75,6 @@
#include "cutlass/gemm/collective/sm120_sparse_mma_tma.hpp"
#include "cutlass/gemm/collective/sm120_blockscaled_sparse_mma_tma.hpp"
#endif // !defined(__CUDACC_RTC__)

#if defined(SYCL_INTEL_TARGET)
#include "cutlass/gemm/collective/xe_mma.hpp"
#include "cutlass/gemm/collective/xe_array_mma.hpp"
#include "cutlass/gemm/collective/xe_array_mma_fp8.hpp"
#include "cutlass/gemm/collective/xe_mma_mixed_input.hpp"
#include "cutlass/gemm/collective/xe_mma_w8a8.hpp"
#include "cutlass/gemm/collective/xe_mma_fp8_scaling.hpp"
#endif

#if defined(CUTLASS_ENABLE_SYCL)
Expand Down
10 changes: 5 additions & 5 deletions include/cutlass/gemm/kernel/gemm_universal.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,6 +53,11 @@ struct IsCutlass3ArrayKernel<ProblemShape, cute::void_t<typename ProblemShape::U

////////////////////////////////////////////////////////////////////////////////

#if defined(SYCL_INTEL_TARGET)
#include "cutlass/gemm/kernel/xe_gemm.hpp"
#include "cutlass/gemm/kernel/xe_gemm_cooperative.hpp"
#include "cutlass/gemm/kernel/xe_gemm_array_cooperative.hpp"
#else
#include "cutlass/gemm/kernel/sm70_gemm.hpp"
#include "cutlass/gemm/kernel/sm90_gemm_tma.hpp"
#include "cutlass/gemm/kernel/sm90_gemm_warpspecialized.hpp"
Expand All @@ -71,11 +76,6 @@ struct IsCutlass3ArrayKernel<ProblemShape, cute::void_t<typename ProblemShape::U
#include "cutlass/gemm/kernel/sm100_gemm_array_tma_warpspecialized_mma_transform.hpp"
#include "cutlass/gemm/kernel/sm100_sparse_gemm_tma_warpspecialized.hpp"
#include "cutlass/gemm/kernel/sm120_gemm_tma_warpspecialized_cooperative_asymmetric_dma.hpp"

#if defined(SYCL_INTEL_TARGET)
#include "cutlass/gemm/kernel/xe_gemm.hpp"
#include "cutlass/gemm/kernel/xe_gemm_cooperative.hpp"
#include "cutlass/gemm/kernel/xe_gemm_array_cooperative.hpp"
#endif

////////////////////////////////////////////////////////////////////////////////
107 changes: 55 additions & 52 deletions include/cutlass/gemm/kernel/tile_scheduler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -60,16 +60,18 @@ struct StaticPersistentScheduler { };
////////////////////////////////////////////////////////////////////////////////

#include "cutlass/gemm/kernel/sm90_tile_scheduler.hpp"
#include "cutlass/gemm/kernel/sm100_static_tile_scheduler.hpp"

#if defined(SYCL_INTEL_TARGET)
#include "cutlass/gemm/kernel/xe_tile_scheduler_streamk.hpp"
#include "cutlass/gemm/kernel/xe_tile_scheduler_group.hpp"
#else
#include "cutlass/gemm/kernel/sm100_static_tile_scheduler.hpp"

#include "cutlass/gemm/kernel/sm90_tile_scheduler_stream_k.hpp"
#include "cutlass/gemm/kernel/sm90_tile_scheduler_group.hpp"
#include "cutlass/gemm/kernel/sm100_tile_scheduler.hpp"
#include "cutlass/gemm/kernel/sm100_tile_scheduler_stream_k.hpp"
#include "cutlass/gemm/kernel/sm100_tile_scheduler_group.hpp"
#if defined (SYCL_INTEL_TARGET)
#include "cutlass/gemm/kernel/xe_tile_scheduler_streamk.hpp"
#include "cutlass/gemm/kernel/xe_tile_scheduler_group.hpp"
#endif
////////////////////////////////////////////////////////////////////////////////

Expand All @@ -96,14 +98,14 @@ template <
class ArchTag,
class TileShape,
class ClusterShape
, uint32_t SchedulerPipelineStageCount
, uint32_t SchedulerPipelineStageCount
>
struct TileSchedulerSelector<
PersistentScheduler,
ArchTag,
TileShape,
ClusterShape
, SchedulerPipelineStageCount
, SchedulerPipelineStageCount
> {
using Scheduler = PersistentTileSchedulerSm90;
};
Expand Down Expand Up @@ -131,6 +133,52 @@ struct TileSchedulerSelector<
>::Scheduler;
};

#if defined(SYCL_INTEL_TARGET)
template <
class TileShape,
class ClusterShape,
uint32_t ThreadsPerBlock
>
struct TileSchedulerSelector<
StreamKScheduler,
arch::IntelXe,
TileShape,
ClusterShape,
ThreadsPerBlock
> {
using Scheduler = PersistentTileSchedulerXeStreamK<TileShape, ThreadsPerBlock>;
};

template <
class TileShape,
class ClusterShape,
uint32_t SchedulerPipelineStageCount,
class GroupProblemShape
>
struct TileSchedulerSelector<
GroupScheduler,
arch::IntelXe,
TileShape,
ClusterShape,
SchedulerPipelineStageCount,
GroupProblemShape
> {
using Scheduler = PersistentTileSchedulerXeGroup<GroupProblemShape>;
};
template <
class TileShape,
class ClusterShape
>
struct TileSchedulerSelector<
PersistentScheduler,
arch::IntelXe,
TileShape,
ClusterShape
> {
using Scheduler = PersistentTileSchedulerSm90;
};
#else

template <
class TileShape,
class ClusterShape
Expand Down Expand Up @@ -179,52 +227,6 @@ struct TileSchedulerSelector<
using Scheduler = PersistentTileSchedulerSm90Group<GroupProblemShape, SchedulerPipelineStageCount>;
};

#if defined (SYCL_INTEL_TARGET)
template <
class TileShape,
class ClusterShape,
uint32_t ThreadsPerBlock
>
struct TileSchedulerSelector<
StreamKScheduler,
arch::IntelXe,
TileShape,
ClusterShape,
ThreadsPerBlock
> {
using Scheduler = PersistentTileSchedulerXeStreamK<TileShape, ThreadsPerBlock>;
};

template <
class TileShape,
class ClusterShape,
uint32_t SchedulerPipelineStageCount,
class GroupProblemShape
>
struct TileSchedulerSelector<
GroupScheduler,
arch::IntelXe,
TileShape,
ClusterShape,
SchedulerPipelineStageCount,
GroupProblemShape
> {
using Scheduler = PersistentTileSchedulerXeGroup<GroupProblemShape>;
};
template <
class TileShape,
class ClusterShape
>
struct TileSchedulerSelector<
PersistentScheduler,
arch::IntelXe,
TileShape,
ClusterShape
> {
using Scheduler = PersistentTileSchedulerSm90;
};
#endif

template <class TileShape, class ClusterShape, uint32_t SchedulerPipelineStageCount>
struct TileSchedulerSelector<
PersistentScheduler,
Expand Down Expand Up @@ -405,6 +407,7 @@ struct TileSchedulerSelector<
> {
using Scheduler = PersistentTileSchedulerSm90Group<GroupProblemShape, SchedulerPipelineStageCount>;
};
#endif

////////////////////////////////////////////////////////////////////////////////

Expand Down
Loading
Loading