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
|
#pragma once
#include <cmath>
#include <cublas_v2.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
#define __fp16 half
__fp16 float2half_rn (float a);
__device__
inline half htanh(const half x)
{
//half ret = ((half)1.0f - hexp((half)-2.0f * x)) / ((half)1.0f + hexp((half)-2.0f * x));
//half ret = (hexp((half)2.0f * x) - (half)1.0f) / (hexp((half)2.0f * x) + (half)1.0f);
//half ret = (hexp(x) - hexp(-x)) / (hexp(x) + hexp(-x));
half ret = tanhf(x);
return ret;
}
namespace thrust
{
namespace detail
{
namespace functional
{
//////////////////////////////////////////////////////////////////////////////////////////
template<typename T>
struct half_unary_logit : public thrust::unary_function<T,T> {
__host__ __device__
T operator()(const T &x) const { return (half)1.0 / ((half)1.0 + hexp(-x)); }
};
template<typename Eval>
__host__ __device__
actor<composite<unary_operator<half_unary_logit>, actor<Eval>>>
HalfLogit(const actor<Eval> &_1) {
return compose(unary_operator<half_unary_logit>(), _1);
}
//////////////////////////////////////////////////////////////////////////////////////////
template<typename T>
struct half_unary_tanh : public thrust::unary_function<T,T> {
__host__ __device__
T operator()(const T &x) const { return htanh(x); }
};
template<typename Eval>
__host__ __device__
actor<composite<unary_operator<half_unary_tanh>, actor<Eval>>>
HalfTanh(const actor<Eval> &_1) {
return compose(unary_operator<half_unary_tanh>(), _1);
}
//////////////////////////////////////////////////////////////////////////////////////////
template<typename T>
struct half_binary_add : public thrust::binary_function<T,T,T> {
__host__ __device__
T operator()(const T &x, const T &y) const { return x + y; }
};
template<typename Eval1, typename Eval2>
__host__ __device__
actor<composite<binary_operator<half_binary_add>, actor<Eval1>, actor<Eval2>>>
HalfAdd(const actor<Eval1> &_1, const actor<Eval2> &_2) {
return compose(binary_operator<half_binary_add>(), _1, _2);
}
//////////////////////////////////////////////////////////////////////////////////////////
}
}
}
|
