forked from pytorch/pytorch
-
Notifications
You must be signed in to change notification settings - Fork 0
/
HardTanh.cu
63 lines (54 loc) · 1.35 KB
/
HardTanh.cu
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
#include "THCUNN.h"
#include "TH/THHalf.h"
#include "THCHalfAutoNumerics.cuh"
#include <THC/THCApply.cuh>
template <typename T>
struct hardtanhupdateOutput_functor
{
const T max_val_;
const T min_val_;
hardtanhupdateOutput_functor(T min_val, T max_val)
: min_val_(min_val)
, max_val_(max_val)
{}
__device__ void operator()(T *output, const T *input) const
{
if (*input < min_val_)
*output = min_val_;
else if (*input > max_val_)
*output = max_val_;
else
*output = *input;
}
__device__ void operator()(T *input) const
{
if (*input < min_val_)
*input = min_val_;
else if (*input > max_val_)
*input = max_val_;
}
};
template <typename T>
struct hardtanhupdateGradInput_functor
{
const T max_val_;
const T min_val_;
hardtanhupdateGradInput_functor(T min_val, T max_val)
: min_val_(min_val)
, max_val_(max_val)
{}
__device__ void operator()(T *gradInput, const T *input, const T *gradOutput) const
{
if (*input <= min_val_ || *input >= max_val_)
*gradInput = ScalarConvert<int, T>::to(0);
else
*gradInput = *gradOutput;
}
__device__ void operator()(T *gradInput, const T *input) const
{
if (*input <= min_val_ || *input >= max_val_)
*gradInput = ScalarConvert<int, T>::to(0);
}
};
#include "generic/HardTanh.cu"
#include "THCGenerateFloatTypes.h"