Skip to content

Commit ec10f2e

Browse files
alalekalalek
committed
Merge pull request #20877 from rogday:simple_layers
2 parents f36c268 + 1feb383 commit ec10f2e

File tree

12 files changed

+936
-703
lines changed

12 files changed

+936
-703
lines changed

‎modules/dnn/include/opencv2/dnn/all_layers.hpp‎

Lines changed: 43 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -600,6 +600,42 @@ CV__DNN_INLINE_NS_BEGIN
600600
static Ptr<ExpLayer> create(const LayerParams &params);
601601
};
602602

603+
class CV_EXPORTS CeilLayer : public ActivationLayer
604+
{
605+
public:
606+
static Ptr<CeilLayer> create(const LayerParams &params);
607+
};
608+
609+
class CV_EXPORTS FloorLayer : public ActivationLayer
610+
{
611+
public:
612+
static Ptr<FloorLayer> create(const LayerParams &params);
613+
};
614+
615+
class CV_EXPORTS LogLayer : public ActivationLayer
616+
{
617+
public:
618+
static Ptr<LogLayer> create(const LayerParams &params);
619+
};
620+
621+
class CV_EXPORTS RoundLayer : public ActivationLayer
622+
{
623+
public:
624+
static Ptr<RoundLayer> create(const LayerParams &params);
625+
};
626+
627+
class CV_EXPORTS SqrtLayer : public ActivationLayer
628+
{
629+
public:
630+
static Ptr<SqrtLayer> create(const LayerParams &params);
631+
};
632+
633+
class CV_EXPORTS NotLayer : public ActivationLayer
634+
{
635+
public:
636+
static Ptr<NotLayer> create(const LayerParams &params);
637+
};
638+
603639
class CV_EXPORTS ActivationLayerInt8 : public ActivationLayer
604640
{
605641
public:
@@ -665,6 +701,7 @@ CV__DNN_INLINE_NS_BEGIN
665701
public:
666702
bool hasBias;
667703
int axis;
704+
String mode;
668705

669706
static Ptr<ScaleLayer> create(const LayerParams& params);
670707
};
@@ -689,6 +726,12 @@ CV__DNN_INLINE_NS_BEGIN
689726
static Ptr<Layer> create(const LayerParams& params);
690727
};
691728

729+
class CV_EXPORTS CompareLayer : public Layer
730+
{
731+
public:
732+
static Ptr<Layer> create(const LayerParams& params);
733+
};
734+
692735
class CV_EXPORTS DataAugmentationLayer : public Layer
693736
{
694737
public:

‎modules/dnn/src/cuda/activations.cu‎

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -128,6 +128,36 @@ void bnll(const Stream& stream, Span<T> output, View<T> input) {
128128
generic_op<T, BNLLFunctor<T>>(stream, output, input);
129129
}
130130

131+
template <class T>
132+
void ceil(const Stream& stream, Span<T> output, View<T> input) {
133+
generic_op<T, CeilFunctor<T>>(stream, output, input);
134+
}
135+
136+
template <class T>
137+
void floor(const Stream& stream, Span<T> output, View<T> input) {
138+
generic_op<T, FloorFunctor<T>>(stream, output, input);
139+
}
140+
141+
template <class T>
142+
void log(const Stream& stream, Span<T> output, View<T> input) {
143+
generic_op<T, LogFunctor<T>>(stream, output, input);
144+
}
145+
146+
template <class T>
147+
void rint(const Stream& stream, Span<T> output, View<T> input) {
148+
generic_op<T, RintFunctor<T>>(stream, output, input);
149+
}
150+
151+
template <class T>
152+
void sqrt(const Stream& stream, Span<T> output, View<T> input) {
153+
generic_op<T, SqrtFunctor<T>>(stream, output, input);
154+
}
155+
156+
template <class T>
157+
void not_k(const Stream& stream, Span<T> output, View<T> input) {
158+
generic_op<T, NotFunctor<T>>(stream, output, input);
159+
}
160+
131161
template <class T>
132162
void abs(const Stream& stream, Span<T> output, View<T> input) {
133163
generic_op<T, AbsFunctor<T>>(stream, output, input);
@@ -160,6 +190,12 @@ template void sigmoid<__half>(const Stream&, Span<__half>, View<__half>);
160190
template void elu<__half>(const Stream&, Span<__half>, View<__half>);
161191
template void abs<__half>(const Stream& stream, Span<__half> output, View<__half> input);
162192
template void bnll<__half>(const Stream&, Span<__half>, View<__half>);
193+
template void ceil<__half>(const Stream&, Span<__half>, View<__half>);
194+
template void floor<__half>(const Stream&, Span<__half>, View<__half>);
195+
template void log<__half>(const Stream&, Span<__half>, View<__half>);
196+
template void rint<__half>(const Stream&, Span<__half>, View<__half>);
197+
template void sqrt<__half>(const Stream&, Span<__half>, View<__half>);
198+
template void not_k<__half>(const Stream&, Span<__half>, View<__half>);
163199
template void power<__half>(const Stream&, Span<__half>, View<__half>, __half, __half, __half);
164200
template void exp<__half>(const Stream&, Span<__half>, View<__half>, __half, __half);
165201
#endif
@@ -174,6 +210,12 @@ template void sigmoid<float>(const Stream&, Span<float>, View<float>);
174210
template void elu<float>(const Stream&, Span<float>, View<float>);
175211
template void abs<float>(const Stream& stream, Span<float> output, View<float> input);
176212
template void bnll<float>(const Stream&, Span<float>, View<float>);
213+
template void ceil<float>(const Stream&, Span<float>, View<float>);
214+
template void floor<float>(const Stream&, Span<float>, View<float>);
215+
template void log<float>(const Stream&, Span<float>, View<float>);
216+
template void rint<float>(const Stream&, Span<float>, View<float>);
217+
template void sqrt<float>(const Stream&, Span<float>, View<float>);
218+
template void not_k<float>(const Stream&, Span<float>, View<float>);
177219
template void power<float>(const Stream&, Span<float>, View<float>, float, float, float);
178220
template void exp<float>(const Stream&, Span<float>, View<float>, float, float);
179221

‎modules/dnn/src/cuda/functors.hpp‎

Lines changed: 90 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -209,6 +209,96 @@ struct BNLLFunctor {
209209
}
210210
};
211211

212+
template <class T>
213+
struct CeilFunctor {
214+
struct Params {
215+
CUDA4DNN_HOST_DEVICE Params() { }
216+
};
217+
218+
CUDA4DNN_DEVICE CeilFunctor() { }
219+
CUDA4DNN_DEVICE CeilFunctor(const Params& params) { }
220+
221+
CUDA4DNN_DEVICE T operator()(T value) {
222+
using csl::device::ceil;
223+
return ceil(value);
224+
}
225+
};
226+
227+
template <class T>
228+
struct FloorFunctor {
229+
struct Params {
230+
CUDA4DNN_HOST_DEVICE Params() { }
231+
};
232+
233+
CUDA4DNN_DEVICE FloorFunctor() { }
234+
CUDA4DNN_DEVICE FloorFunctor(const Params& params) { }
235+
236+
CUDA4DNN_DEVICE T operator()(T value) {
237+
using csl::device::floor;
238+
return floor(value);
239+
}
240+
};
241+
242+
template <class T>
243+
struct LogFunctor {
244+
struct Params {
245+
CUDA4DNN_HOST_DEVICE Params() { }
246+
};
247+
248+
CUDA4DNN_DEVICE LogFunctor() { }
249+
CUDA4DNN_DEVICE LogFunctor(const Params& params) { }
250+
251+
CUDA4DNN_DEVICE T operator()(T value) {
252+
using csl::device::log;
253+
return log(value);
254+
}
255+
};
256+
257+
template <class T>
258+
struct RintFunctor {
259+
struct Params {
260+
CUDA4DNN_HOST_DEVICE Params() { }
261+
};
262+
263+
CUDA4DNN_DEVICE RintFunctor() { }
264+
CUDA4DNN_DEVICE RintFunctor(const Params& params) { }
265+
266+
CUDA4DNN_DEVICE T operator()(T value) {
267+
using csl::device::rint;
268+
return rint(value);
269+
}
270+
};
271+
272+
template <class T>
273+
struct SqrtFunctor {
274+
struct Params {
275+
CUDA4DNN_HOST_DEVICE Params() { }
276+
};
277+
278+
CUDA4DNN_DEVICE SqrtFunctor() { }
279+
CUDA4DNN_DEVICE SqrtFunctor(const Params& params) { }
280+
281+
CUDA4DNN_DEVICE T operator()(T value) {
282+
using csl::device::sqrt;
283+
return sqrt(value);
284+
}
285+
};
286+
287+
template <class T>
288+
struct NotFunctor {
289+
struct Params {
290+
CUDA4DNN_HOST_DEVICE Params() { }
291+
};
292+
293+
CUDA4DNN_DEVICE NotFunctor() { }
294+
CUDA4DNN_DEVICE NotFunctor(const Params& params) { }
295+
296+
CUDA4DNN_DEVICE T operator()(T value) {
297+
using csl::device::floor;
298+
return floor(static_cast<T>(1.) - value);
299+
}
300+
};
301+
212302
template <class T>
213303
struct PowerFunctor {
214304
struct Params {

‎modules/dnn/src/cuda/math.hpp‎

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -119,6 +119,27 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace csl { namespace de
119119
template <> inline __device__ __half round(__half value) { return hrint(value); }
120120
#endif
121121

122+
template <class T> __device__ T floor(T value);
123+
template <> inline __device__ double floor(double value) { return ::floor(value); }
124+
template <> inline __device__ float floor(float value) { return floorf(value); }
125+
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
126+
template <> inline __device__ __half floor(__half value) { return hfloor(value); }
127+
#endif
128+
129+
template <class T> __device__ T log(T value);
130+
template <> inline __device__ double log(double value) { return ::log(value); }
131+
template <> inline __device__ float log(float value) { return logf(value); }
132+
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
133+
template <> inline __device__ __half log(__half value) { return hlog(value); }
134+
#endif
135+
136+
template <class T> __device__ T rint(T value);
137+
template <> inline __device__ double rint(double value) { return ::rint(value); }
138+
template <> inline __device__ float rint(float value) { return rintf(value); }
139+
#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 530)
140+
template <> inline __device__ __half rint(__half value) { return hrint(value); }
141+
#endif
142+
122143
template <class T> __device__ T ceil(T value);
123144
template <> inline __device__ double ceil(double value) { return ::ceil(value); }
124145
template <> inline __device__ float ceil(float value) { return ceilf(value); }

‎modules/dnn/src/cuda4dnn/kernels/activations.hpp‎

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -42,6 +42,24 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels {
4242
template <class T>
4343
void bnll(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input);
4444

45+
template <class T>
46+
void ceil(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input);
47+
48+
template <class T>
49+
void floor(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input);
50+
51+
template <class T>
52+
void log(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input);
53+
54+
template <class T>
55+
void rint(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input);
56+
57+
template <class T>
58+
void sqrt(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input);
59+
60+
template <class T>
61+
void not_k(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input);
62+
4563
template <class T>
4664
void power(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input, T exp, T scale, T shift);
4765

0 commit comments

Comments
 (0)