Skip to content

GitLab

Switch branch/tag
Find file Normal viewHistoryPermalink
static_switch.h 5.1 KB
Newer Older
Tri Dao's avatar
FA3 initial code release  
Tri Dao committed
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83 
// Inspired by
// https://github.com/NVIDIA/DALI/blob/main/include/dali/core/static_switch.h
// and https://github.com/pytorch/pytorch/blob/master/aten/src/ATen/Dispatch.h

#pragma once

/// @param COND       - a boolean expression to switch by
/// @param CONST_NAME - a name given for the constexpr bool variable.
/// @param ...       - code to execute for true and false
///
/// Usage:
/// ```
/// BOOL_SWITCH(flag, BoolConst, [&] {
///     some_function<BoolConst>(...);
/// });
/// ```
//

#define BOOL_SWITCH(COND, CONST_NAME, ...)                                     \
  [&] {                                                                        \
    if (COND) {                                                                \
      constexpr static bool CONST_NAME = true;                                 \
      return __VA_ARGS__();                                                    \
    } else {                                                                   \
      constexpr static bool CONST_NAME = false;                                \
      return __VA_ARGS__();                                                    \
    }                                                                          \
  }()

#define PREC_SWITCH(PRECTYPE, ...)                                             \
  [&] {                                                                        \
    if (PRECTYPE == 1) {                                                       \
      using kPrecType = cutlass::half_t;                                       \
      constexpr static bool kSoftFp16 = false;                                 \
      constexpr static bool kHybrid = false;                                   \
      return __VA_ARGS__();                                                    \
    } else if (PRECTYPE == 2) {                                                \
      using kPrecType = cutlass::float_e4m3_t;                                 \
      constexpr static bool kSoftFp16 = false;                                 \
      constexpr static bool kHybrid = false;                                   \
      return __VA_ARGS__();                                                    \
    } else if (PRECTYPE == 3) {                                                \
      using kPrecType = cutlass::float_e4m3_t;                                 \
      constexpr static bool kSoftFp16 = false;                                 \
      constexpr static bool kHybrid = true;                                    \
      return __VA_ARGS__();                                                    \
    } else if (PRECTYPE == 4) {                                                \
      using kPrecType = cutlass::float_e4m3_t;                                 \
      constexpr static bool kSoftFp16 = true;                                  \
      constexpr static bool kHybrid = false;                                   \
      return __VA_ARGS__();                                                    \
    }                                                                          \
  }()

#define HEADDIM_SWITCH(HEADDIM, ...)                                           \
  [&] {                                                                        \
    if (HEADDIM == 64) {                                                       \
      constexpr static int kHeadSize = 64;                                     \
      return __VA_ARGS__();                                                    \
    } else if (HEADDIM == 128) {                                               \
      constexpr static int kHeadSize = 128;                                    \
      return __VA_ARGS__();                                                    \
    } else if (HEADDIM == 256) {                                               \
      constexpr static int kHeadSize = 256;                                    \
      return __VA_ARGS__();                                                    \
    }                                                                          \
  }()

#define SEQLEN_SWITCH(USE_VAR_SEQ_LEN, SEQ_LEN_OUT_OF_BOUND_CHECK, ...)        \
  [&] {                                                                        \
    if (!USE_VAR_SEQ_LEN) {                                                    \
      if (SEQ_LEN_OUT_OF_BOUND_CHECK) {                                        \
        using kSeqLenTraitsType = FixedSeqLenTraits<true>;                     \
        return __VA_ARGS__();                                                  \
      } else {                                                                 \
        using kSeqLenTraitsType = FixedSeqLenTraits<false>;                    \
        return __VA_ARGS__();                                                  \
      }                                                                        \
    } else {                                                                   \
      using kSeqLenTraitsType = VarSeqLenTraits;                               \
      return __VA_ARGS__();                                                    \
    }                                                                          \
  }()