static_switch.h 1.65 KB
Newer Older
Tri Dao's avatar
Tri Dao committed
1
2
// 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
3
// and https://github.com/facebookresearch/xformers/blob/main/xformers/csrc/attention/cuda/fmha/gemm_kernel_utils.h#L8
Tri Dao's avatar
Tri Dao committed
4
5
6
7
8
9
10
11
12

#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:
/// ```
13
/// BOOL_SWITCH(flag, BoolConst, ([&] {
Tri Dao's avatar
Tri Dao committed
14
///     some_function<BoolConst>(...);
15
/// }));
Tri Dao's avatar
Tri Dao committed
16
/// ```
17
/// We need "({" and "})" to make sure that the code is a single argument being passed to the macro.
18
19
20
21
22
23
24
25
26
#define BOOL_SWITCH(COND, CONST_NAME, F)       \
    {                                          \
        if (COND) {                            \
            constexpr bool CONST_NAME = true;  \
            F();                               \
        } else {                               \
            constexpr bool CONST_NAME = false; \
            F();                               \
        }                                      \
27
28
29
30
    }

// modified from BOOL_SWITCH
// because MSVC cannot handle std::conditional with constexpr variable
31
32
33
34
35
36
37
38
39
40
#define FP16_SWITCH(COND, F)                 \
    {                                        \
        if (COND) {                          \
            using elem_type = __nv_bfloat16; \
            F();                             \
        } else {                             \
            using elem_type = __half;        \
            F();                             \
        }                                    \
    }