"src/turbomind/kernels/sampling_topp_kernels.h" did not exist on "720fc533da804ac3f46ee938864403e51fcd9fa7"
reduction_functions_accumulate.hpp 4.66 KB
Newer Older
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
/*******************************************************************************
 *
 * MIT License
 *
 * Copyright (c) 2020 Advanced Micro Devices, Inc.
 *
 * Permission is hereby granted, free of charge, to any person obtaining a copy
 * of this software and associated documentation files (the "Software"), to deal
 * in the Software without restriction, including without limitation the rights
 * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
 * copies of the Software, and to permit persons to whom the Software is
 * furnished to do so, subject to the following conditions:
 *
 * The above copyright notice and this permission notice shall be included in all
 * copies or substantial portions of the Software.
 *
 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
 * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
 * SOFTWARE.
 *
 *******************************************************************************/
#ifndef CK_REDUCTION_FUNCTIONS_BINOP_HPP
#define CK_REDUCTION_FUNCTIONS_BINOP_HPP

#include "data_type.hpp"
30
#include "math_v2.hpp"
31
32
33
34
35
36
37

#include "reduction_common.hpp"
#include "reduction_operator.hpp"

namespace ck {
namespace detail {

38
39
40
41
42
43
44
45
46
47
48
49
50
// Check for NaN; guarantee NaNs are NOT propagated to result (i.e., ignore NaNs)
template <typename ReduceOperation, typename AccDataType>
struct AccumulateWithNanIgnore
{
    __device__ static inline void Calculate(AccDataType& accuVal, AccDataType currVal)
    {
        if(!isnan(currVal))
        {
            ReduceOperation{}(accuVal, currVal);
        }
    };
};

51
52
template <bool PropagateNan, typename ReduceOperation, typename AccDataType>
struct AccumulateWithNanCheck;
53

54
55
56
57
58
// Does not check for NaN; does not guarantee NaNs be propagated to result
// e.g., given that max(a, b) = a > b ? a : b
// then  max(NaN, 1) returns 1
//       max(1, NaN) returns NaN
// since any comparison involving NaNs returns false
59
60
template <typename ReduceOperation, typename AccDataType>
struct AccumulateWithNanCheck<false, ReduceOperation, AccDataType>
61
62
{
    // cppcheck-suppress constParameter
63
    __host__ __device__ static inline void Calculate(AccDataType& accuVal, AccDataType currVal)
64
65
66
67
68
    {
        ReduceOperation{}(accuVal, currVal);
    };
};

69
// Check for NaN; guarantees NaNs be propagated to result
70
71
72
template <typename ReduceOperation, typename AccDataType>
struct AccumulateWithNanCheck<true, ReduceOperation, AccDataType>
{
73
    __host__ __device__ static inline void Calculate(AccDataType& accuVal, AccDataType currVal)
74
    {
75
76
77
        using ck::math::isnan;

        if(isnan(currVal))
78
79
80
81
82
83
84
        {
            accuVal = currVal;
        }
        else
        {
            ReduceOperation{}(accuVal, currVal);
        };
85
    };
86
87
88
89
};

template <bool PropagateNan, typename ReduceOperation, typename AccDataType, typename IndexDataType>
struct AccumulateWithIndexAndNanCheck;
90

91
92
93
template <typename ReduceOperation, typename AccDataType, typename IndexDataType>
struct AccumulateWithIndexAndNanCheck<false, ReduceOperation, AccDataType, IndexDataType>
{
94
    __host__ __device__ static inline void
95
    // cppcheck-suppress constParameter
96
97
98
99
    Calculate(AccDataType& accuVal,
              AccDataType currVal,
              IndexDataType& accuIndex,
              IndexDataType currIndex)
100
101
102
    {
        bool changed = false;

103
        ReduceOperation{}(accuVal, currVal, changed);
104
105
106
107
108
109

        if(changed)
            accuIndex = currIndex;
    };
};

110
111
template <typename ReduceOperation, typename AccDataType, typename IndexDataType>
struct AccumulateWithIndexAndNanCheck<true, ReduceOperation, AccDataType, IndexDataType>
112
{
113
    // The method is called when the ReduceOperation is indexable and the user asked for indices
114
115
116
117
    __host__ __device__ static inline void Calculate(AccDataType& accuVal,
                                                     AccDataType currVal,
                                                     IndexDataType& accuIndex,
                                                     IndexDataType currIndex)
118
    {
119
120
121
        using ck::math::isnan;

        if(isnan(currVal))
122
123
124
125
126
127
128
129
        {
            accuVal   = currVal;
            accuIndex = currIndex;
        }
        else
        {
            bool changed = false;

130
            ReduceOperation{}(accuVal, currVal, changed);
131
132
133
134
135
136
137
138
139
140
141

            if(changed)
                accuIndex = currIndex;
        }
    };
};

}; // namespace detail
}; // end of namespace ck

#endif