Logo ROOT  
Reference Guide
 
Loading...
Searching...
No Matches
ActivationFunctions.cu
Go to the documentation of this file.
1// @(#)root/tmva/tmva/dnn:$Id$
2// Author: Simon Pfreundschuh 13/07/16
3
4/*************************************************************************
5 * Copyright (C) 2016, Simon Pfreundschuh *
6 * All rights reserved. *
7 * *
8 * For the licensing terms see $ROOTSYS/LICENSE. *
9 * For the list of contributors see $ROOTSYS/README/CREDITS. *
10 *************************************************************************/
11
12 //////////////////////////////////////////////////////////////////
13 // Implementation of the activation functions for the TCuda //
14 // implementation of the low-level interface. //
15 //////////////////////////////////////////////////////////////////
16
18/*#include "TMVA/DNN/Architectures/Cuda.h"
19#include "TMVA/DNN/Architectures/Cuda/Device.h"
20#include "Kernels.cuh"*/
21
22namespace TMVA
23{
24namespace DNN
25{
26
27//______________________________________________________________________________
28template<typename AFloat>
29void TCudnn<AFloat>::ActivationFunctionForward(Tensor_t & X, EActivationFunction activFunct, const ActivationDescriptor_t activationDescr, const double coef, const AFloat alpha, const AFloat beta)
30{
31 // compute forward activation in place
32 // Nothing to do for identity function
33 if (activFunct == EActivationFunction::kIdentity) return;
34
35 CUDNNCHECK(cudnnActivationForward(X.GetCudnnHandle(),
36 activationDescr,
37 &alpha,
38 X.GetTensorDescriptor(),
39 X.GetDataPointer(),
40 &beta,
41 X.GetTensorDescriptor(), // Can be computed in place
42 X.GetDataPointer()));
43}
44//______________________________________________________________________________
45template <typename AFloat>
46void TCudnn<AFloat>::ActivationFunctionForward(Tensor_t & Y, const Tensor_t &X, EActivationFunction activFunct,
47 const ActivationDescriptor_t activationDescr, const double coef,
48 const AFloat alpha, const AFloat beta)
49{
50 // compute forward activation with different input/output tensor (needed in training)
51 // Nothing to do for identity function
52 if (activFunct == EActivationFunction::kIdentity) {
53 TCudnn<AFloat>::Copy(Y, X);
54 return;
55 }
56
57 CUDNNCHECK(cudnnActivationForward(X.GetCudnnHandle(), activationDescr, &alpha, X.GetTensorDescriptor(),
58 X.GetDataPointer(), &beta,
59 Y.GetTensorDescriptor(), // Can be computed in place
60 Y.GetDataPointer()));
61}
62
63//______________________________________________________________________________
64template<typename AFloat>
65void TCudnn<AFloat>::ActivationFunctionBackward(Tensor_t & dX, const Tensor_t & Y,
66 const Tensor_t & dY, const Tensor_t & X,
67 EActivationFunction activFunct ,
68 const ActivationDescriptor_t activationDescr,
69 const AFloat alpha, const AFloat beta)
70{
71 // For identity function output dX is = dY
72 if (activFunct == EActivationFunction::kIdentity) {
73 Copy(dX,dY);
74 return;
75 }
76 //std::cout << "No identityy\n";
77 //Y.Print();
78 // The activation descriptor is set in the forward pass
79 CUDNNCHECK(cudnnActivationBackward(X.GetCudnnHandle(),
80 activationDescr,
81 &alpha,
82 Y.GetTensorDescriptor(),
83 Y.GetDataPointer(),
84 dY.GetTensorDescriptor(),
85 dY.GetDataPointer(),
86 X.GetTensorDescriptor(),
87 X.GetDataPointer(),
88 &beta,
89 dX.GetTensorDescriptor(),
90 dX.GetDataPointer()));
91}
92
93#if 0
94//______________________________________________________________________________
95/*template<typename AFloat>
96void TCudnn<AFloat>::Relu(Tensor_t & X, ActivationDescriptor_t activationDescr, const double coef, const AFloat alpha, const AFloat beta)
97{
98 Activation(X, EActivationFunction::kRelu, activationDescr, coef, alpha, beta);
99}
100
101//______________________________________________________________________________
102template<typename AFloat>
103void TCudnn<AFloat>::ReluDerivative(const Tensor_t & Y, const Tensor_t & dY,
104 const Tensor_t & X, Tensor_t & dX,
105 const ActivationDescriptor_t activationDescr,
106 const AFloat alpha, const AFloat beta)
107{
108 ActivationFunctionBackward(Y, dY, X, dX, activationDescr, alpha, beta);
109}
110
111//______________________________________________________________________________
112template<typename AFloat>
113void TCudnn<AFloat>::Sigmoid(Tensor_t & X, ActivationDescriptor_t activationDescr, const double coef, const AFloat alpha, const AFloat beta)
114{
115 Activation(X, EActivationFunction::kSigmoid, activationDescr, coef, alpha, beta);
116}
117
118//______________________________________________________________________________
119template<typename AFloat>
120void TCudnn<AFloat>::SigmoidDerivative(const Tensor_t & Y, const Tensor_t & dY,
121 const Tensor_t & X, Tensor_t & dX,
122 const ActivationDescriptor_t activationDescr,
123 const AFloat alpha, const AFloat beta)
124{
125 ActivationFunctionBackward(Y, dY, X, dX, activationDescr, alpha, beta);
126}
127
128//______________________________________________________________________________
129template<typename AFloat>
130void TCudnn<AFloat>::Tanh(Tensor_t & X, ActivationDescriptor_t activationDescr, const double coef, const AFloat alpha, const AFloat beta)
131{
132 Activation(X, EActivationFunction::kTanh, activationDescr, alpha, beta);
133}
134
135//______________________________________________________________________________
136template<typename AFloat>
137void TCudnn<AFloat>::TanhDerivative(const Tensor_t & Y, const Tensor_t & dY,
138 const Tensor_t & X, Tensor_t & dX,
139 const ActivationDescriptor_t activationDescr,
140 const AFloat alpha, const AFloat beta)
141{
142 ActivationFunctionBackward(Y, dY, X, dX, activationDescr, alpha, beta);
143}*/
144
145//______________________________________________________________________________
146/*template<typename AFloat>
147void TCudnn<AFloat>::SymmetricRelu(Tensor_t & A)
148{
149 dim3 blockDims = TDevice::BlockDims2D();
150 dim3 gridDims = TDevice::GridDims2D(A);
151 cudaStream_t s = A.GetComputeStream();
152 ::TMVA::DNN::Cuda::SymmetricRelu<<<gridDims, blockDims, 0, s>>>(
153 A.GetDataPointer(),
154 (int) A.GetNrows(),
155 (int) A.GetNcols());
156}
157
158//______________________________________________________________________________
159template<typename AFloat>
160void TCudnn<AFloat>::SymmetricReluDerivative(Tensor_t & B,
161 const Tensor_t & A)
162{
163 dim3 blockDims = TDevice::BlockDims2D();
164 dim3 gridDims = TDevice::GridDims2D(B);
165 cudaStream_t s = A.GetComputeStream();
166 ::TMVA::DNN::Cuda::SymmetricReluDerivative<<<gridDims, blockDims, 0, s>>>(
167 B.GetDataPointer(),
168 A.GetDataPointer(),
169 (int) A.GetNrows(),
170 (int) A.GetNcols());
171 B.SetComputeStream(s);
172}
173
174//______________________________________________________________________________
175template<typename AFloat>
176void TCudnn<AFloat>::SoftSign(Tensor_t & A)
177{
178 dim3 blockDims = TDevice::BlockDims2D();
179 dim3 gridDims = TDevice::GridDims2D(A);
180 cudaStream_t s = A.GetComputeStream();
181 ::TMVA::DNN::Cuda::SoftSign<<<gridDims, blockDims, 0, s>>>(
182 A.GetDataPointer(),
183 (int) A.GetNrows(),
184 (int) A.GetNcols());
185}
186
187//______________________________________________________________________________
188template<typename AFloat>
189void TCudnn<AFloat>::SoftSignDerivative(Tensor_t & B,
190 const Tensor_t & A)
191{
192 dim3 blockDims = TDevice::BlockDims2D();
193 dim3 gridDims = TDevice::GridDims2D(B);
194 cudaStream_t s = A.GetComputeStream();
195 ::TMVA::DNN::Cuda::SoftSignDerivative<<<gridDims, blockDims, 0, s>>>(
196 B.GetDataPointer(),
197 A.GetDataPointer(),
198 (int) A.GetNrows(),
199 (int) A.GetNcols());
200 B.SetComputeStream(s);
201}
202
203//______________________________________________________________________________
204template<typename AFloat>
205void TCudnn<AFloat>::Gauss(Tensor_t & A)
206{
207 dim3 blockDims = TDevice::BlockDims2D();
208 dim3 gridDims = TDevice::GridDims2D(A);
209 cudaStream_t s = A.GetComputeStream();
210 ::TMVA::DNN::Cuda::Gauss<<<gridDims, blockDims, 0, s>>>(
211 A.GetDataPointer(),
212 (int) A.GetNrows(),
213 (int) A.GetNcols());
214}
215
216//______________________________________________________________________________
217template<typename AFloat>
218void TCudnn<AFloat>::GaussDerivative(Tensor_t & B,
219 const Tensor_t & A)
220{
221 dim3 blockDims = TDevice::BlockDims2D();
222 dim3 gridDims = TDevice::GridDims2D(B);
223 cudaStream_t s = A.GetComputeStream();
224 ::TMVA::DNN::Cuda::GaussDerivative<<<gridDims, blockDims, 0, s>>>(
225 B.GetDataPointer(),
226 A.GetDataPointer(),
227 (int) A.GetNrows(),
228 (int) A.GetNcols());
229 B.SetComputeStream(s);
230}*/
231#endif
232
233} // namespace DNN
234} // namespace TMVA
void Copy(void *source, void *dest)
EActivationFunction
Enum that represents layer activation functions.
Definition Functions.h:32
create variable transformations