/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_scheduler.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_scheduler.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-6.4.3/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_scheduler.hpp Source File
gemm_pipeline_ag_bg_cr_scheduler.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #include <ostream>
7 
8 #include "ck_tile/core.hpp"
9 
10 namespace ck_tile {
11 
13 {
14  Default,
15  Intrawave,
16  Interwave,
17 };
18 
19 enum struct TailNumber
20 {
21  // Single / Double buffer pipeline
22  Odd,
23  Even,
24 
25  // Long prefetch pipeline, up to 8
26  One,
27  Two,
28  Three,
29  Four,
30  Five,
31  Six,
32  Seven,
33 
34  // Unroll stages > Prefetch stages, number of loop is multiple of unroll stages
35  Empty,
36  // Unroll stages <= Prefetch stages, number of loop is multiple of unroll stages add
37  // prefetchstages
38  Full,
39 };
40 
41 } // namespace ck_tile
42 
43 inline std::ostream& operator<<(std::ostream& os, const ck_tile::GemmPipelineScheduler& s)
44 {
45  switch(s)
46  {
47  case ck_tile::GemmPipelineScheduler::Default: os << "Default"; break;
48  case ck_tile::GemmPipelineScheduler::Intrawave: os << "Intrawave"; break;
49  case ck_tile::GemmPipelineScheduler::Interwave: os << "Interwave"; break;
50  default: os << "";
51  }
52  return os;
53 }
54 
55 inline std::ostream& operator<<(std::ostream& os, const ck_tile::TailNumber& s)
56 {
57  switch(s)
58  {
59  case ck_tile::TailNumber::Odd: os << "Odd"; break;
60  case ck_tile::TailNumber::Even: os << "Even"; break;
61  case ck_tile::TailNumber::One: os << "One"; break;
62  case ck_tile::TailNumber::Two: os << "Two"; break;
63  case ck_tile::TailNumber::Three: os << "Three"; break;
64  case ck_tile::TailNumber::Four: os << "Four"; break;
65  case ck_tile::TailNumber::Five: os << "Five"; break;
66  case ck_tile::TailNumber::Six: os << "Six"; break;
67  case ck_tile::TailNumber::Seven: os << "Seven"; break;
68  case ck_tile::TailNumber::Empty: os << "Empty"; break;
69  case ck_tile::TailNumber::Full: os << "Full"; break;
70  default: os << "";
71  }
72  return os;
73 }
std::ostream & operator<<(std::ostream &os, const ck_tile::GemmPipelineScheduler &s)
Definition: gemm_pipeline_ag_bg_cr_scheduler.hpp:43
Definition: cluster_descriptor.hpp:13
TailNumber
Definition: gemm_pipeline_ag_bg_cr_scheduler.hpp:20
GemmPipelineScheduler
Definition: gemm_pipeline_ag_bg_cr_scheduler.hpp:13