Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
Menu
Open sidebar
gaoqiong
composable_kernel
Commits
fc870c69
Commit
fc870c69
authored
Oct 18, 2023
by
Bartlomiej Kocot
Browse files
Fixes
parent
3aa0214f
Changes
12
Hide whitespace changes
Inline
Side-by-side
Showing
12 changed files
with
32 additions
and
32 deletions
+32
-32
example/61_conv_fwd_activ/convnd_fwd_activ_common.hpp
example/61_conv_fwd_activ/convnd_fwd_activ_common.hpp
+0
-8
example/61_conv_fwd_activ/convnd_fwd_xdl_abs_fp16.cpp
example/61_conv_fwd_activ/convnd_fwd_xdl_abs_fp16.cpp
+1
-1
example/61_conv_fwd_activ/convnd_fwd_xdl_clippedrelu_fp16.cpp
...ple/61_conv_fwd_activ/convnd_fwd_xdl_clippedrelu_fp16.cpp
+1
-1
example/61_conv_fwd_activ/convnd_fwd_xdl_elu_fp16.cpp
example/61_conv_fwd_activ/convnd_fwd_xdl_elu_fp16.cpp
+1
-1
example/61_conv_fwd_activ/convnd_fwd_xdl_leakyrelu_fp16.cpp
example/61_conv_fwd_activ/convnd_fwd_xdl_leakyrelu_fp16.cpp
+1
-1
example/61_conv_fwd_activ/convnd_fwd_xdl_pow_fp16.cpp
example/61_conv_fwd_activ/convnd_fwd_xdl_pow_fp16.cpp
+1
-1
example/61_conv_fwd_activ/convnd_fwd_xdl_relu_fp16.cpp
example/61_conv_fwd_activ/convnd_fwd_xdl_relu_fp16.cpp
+1
-1
example/61_conv_fwd_activ/convnd_fwd_xdl_sigmoid_fp16.cpp
example/61_conv_fwd_activ/convnd_fwd_xdl_sigmoid_fp16.cpp
+1
-1
example/61_conv_fwd_activ/convnd_fwd_xdl_softrelu_fp16.cpp
example/61_conv_fwd_activ/convnd_fwd_xdl_softrelu_fp16.cpp
+1
-1
example/61_conv_fwd_activ/convnd_fwd_xdl_tanh_fp16.cpp
example/61_conv_fwd_activ/convnd_fwd_xdl_tanh_fp16.cpp
+1
-1
example/61_conv_fwd_activ/run_convnd_fwd_activ_example.inc
example/61_conv_fwd_activ/run_convnd_fwd_activ_example.inc
+9
-1
include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp
...or_operation/gpu/element/unary_element_wise_operation.hpp
+14
-14
No files found.
example/61_conv_fwd_activ/convnd_fwd_activ_common.hpp
View file @
fc870c69
...
...
@@ -94,14 +94,6 @@ using DeviceGroupedConvNDFwdInstance =
S
<
1
,
32
,
1
,
8
>
,
8
>
;
void
print_helper_msg
()
{
std
::
cout
<<
"arg1: verification (0=no, 1=yes)
\n
"
<<
"arg2: initialization (0=no init, 1=integer value, 2=decimal value)
\n
"
<<
"arg3: time kernel (0=no, 1=yes)
\n
"
<<
ck
::
utils
::
conv
::
get_conv_param_parser_helper_msg
()
<<
std
::
endl
;
}
template
<
ck
::
index_t
NDimSpatial
,
typename
InDataType
,
typename
WeiDataType
,
...
...
example/61_conv_fwd_activ/convnd_fwd_xdl_abs_fp16.cpp
View file @
fc870c69
...
...
@@ -8,4 +8,4 @@ using OutElementOp = ck::tensor_operation::element_wise::UnaryAbs;
using
DeviceGroupedConvNDFwdActivInstance
=
DeviceGroupedConvNDFwdInstance
<
OutElementOp
>
;
#include "run_convnd_fwd_activ_example.inc"
int
main
(
int
argc
,
char
*
argv
[])
{
return
run_convnd_fwd_example
(
argc
,
argv
)
?
0
:
1
;
}
int
main
(
int
argc
,
char
*
argv
[])
{
return
!
run_convnd_fwd_example
(
argc
,
argv
);
}
example/61_conv_fwd_activ/convnd_fwd_xdl_clippedrelu_fp16.cpp
View file @
fc870c69
...
...
@@ -8,4 +8,4 @@ using OutElementOp = ck::tensor_operation::element_wise::ClippedRelu;
using
DeviceGroupedConvNDFwdActivInstance
=
DeviceGroupedConvNDFwdInstance
<
OutElementOp
>
;
#include "run_convnd_fwd_activ_example.inc"
int
main
(
int
argc
,
char
*
argv
[])
{
return
run_convnd_fwd_example
(
argc
,
argv
)
?
0
:
1
;
}
int
main
(
int
argc
,
char
*
argv
[])
{
return
!
run_convnd_fwd_example
(
argc
,
argv
);
}
example/61_conv_fwd_activ/convnd_fwd_xdl_elu_fp16.cpp
View file @
fc870c69
...
...
@@ -8,4 +8,4 @@ using OutElementOp = ck::tensor_operation::element_wise::Elu;
using
DeviceGroupedConvNDFwdActivInstance
=
DeviceGroupedConvNDFwdInstance
<
OutElementOp
>
;
#include "run_convnd_fwd_activ_example.inc"
int
main
(
int
argc
,
char
*
argv
[])
{
return
run_convnd_fwd_example
(
argc
,
argv
)
?
0
:
1
;
}
int
main
(
int
argc
,
char
*
argv
[])
{
return
!
run_convnd_fwd_example
(
argc
,
argv
);
}
example/61_conv_fwd_activ/convnd_fwd_xdl_leakyrelu_fp16.cpp
View file @
fc870c69
...
...
@@ -8,4 +8,4 @@ using OutElementOp = ck::tensor_operation::element_wise::LeakyRelu;
using
DeviceGroupedConvNDFwdActivInstance
=
DeviceGroupedConvNDFwdInstance
<
OutElementOp
>
;
#include "run_convnd_fwd_activ_example.inc"
int
main
(
int
argc
,
char
*
argv
[])
{
return
run_convnd_fwd_example
(
argc
,
argv
)
?
0
:
1
;
}
int
main
(
int
argc
,
char
*
argv
[])
{
return
!
run_convnd_fwd_example
(
argc
,
argv
);
}
example/61_conv_fwd_activ/convnd_fwd_xdl_pow_fp16.cpp
View file @
fc870c69
...
...
@@ -8,4 +8,4 @@ using OutElementOp = ck::tensor_operation::element_wise::Power;
using
DeviceGroupedConvNDFwdActivInstance
=
DeviceGroupedConvNDFwdInstance
<
OutElementOp
>
;
#include "run_convnd_fwd_activ_example.inc"
int
main
(
int
argc
,
char
*
argv
[])
{
return
run_convnd_fwd_example
(
argc
,
argv
)
?
0
:
1
;
}
int
main
(
int
argc
,
char
*
argv
[])
{
return
!
run_convnd_fwd_example
(
argc
,
argv
);
}
example/61_conv_fwd_activ/convnd_fwd_xdl_relu_fp16.cpp
View file @
fc870c69
...
...
@@ -8,4 +8,4 @@ using OutElementOp = ck::tensor_operation::element_wise::Relu;
using
DeviceGroupedConvNDFwdActivInstance
=
DeviceGroupedConvNDFwdInstance
<
OutElementOp
>
;
#include "run_convnd_fwd_activ_example.inc"
int
main
(
int
argc
,
char
*
argv
[])
{
return
run_convnd_fwd_example
(
argc
,
argv
)
?
0
:
1
;
}
int
main
(
int
argc
,
char
*
argv
[])
{
return
!
run_convnd_fwd_example
(
argc
,
argv
);
}
example/61_conv_fwd_activ/convnd_fwd_xdl_sigmoid_fp16.cpp
View file @
fc870c69
...
...
@@ -8,4 +8,4 @@ using OutElementOp = ck::tensor_operation::element_wise::Sigmoid;
using
DeviceGroupedConvNDFwdActivInstance
=
DeviceGroupedConvNDFwdInstance
<
OutElementOp
>
;
#include "run_convnd_fwd_activ_example.inc"
int
main
(
int
argc
,
char
*
argv
[])
{
return
run_convnd_fwd_example
(
argc
,
argv
)
?
0
:
1
;
}
int
main
(
int
argc
,
char
*
argv
[])
{
return
!
run_convnd_fwd_example
(
argc
,
argv
);
}
example/61_conv_fwd_activ/convnd_fwd_xdl_softrelu_fp16.cpp
View file @
fc870c69
...
...
@@ -8,4 +8,4 @@ using OutElementOp = ck::tensor_operation::element_wise::SoftRelu;
using
DeviceGroupedConvNDFwdActivInstance
=
DeviceGroupedConvNDFwdInstance
<
OutElementOp
>
;
#include "run_convnd_fwd_activ_example.inc"
int
main
(
int
argc
,
char
*
argv
[])
{
return
run_convnd_fwd_example
(
argc
,
argv
)
?
0
:
1
;
}
int
main
(
int
argc
,
char
*
argv
[])
{
return
!
run_convnd_fwd_example
(
argc
,
argv
);
}
example/61_conv_fwd_activ/convnd_fwd_xdl_tanh_fp16.cpp
View file @
fc870c69
...
...
@@ -8,4 +8,4 @@ using OutElementOp = ck::tensor_operation::element_wise::TanH;
using
DeviceGroupedConvNDFwdActivInstance
=
DeviceGroupedConvNDFwdInstance
<
OutElementOp
>
;
#include "run_convnd_fwd_activ_example.inc"
int
main
(
int
argc
,
char
*
argv
[])
{
return
run_convnd_fwd_example
(
argc
,
argv
)
?
0
:
1
;
}
int
main
(
int
argc
,
char
*
argv
[])
{
return
!
run_convnd_fwd_example
(
argc
,
argv
);
}
example/61_conv_fwd_activ/run_convnd_fwd_activ_example.inc
View file @
fc870c69
...
...
@@ -3,6 +3,14 @@
#pragma once
void
print_helper_msg
()
{
std
::
cout
<<
"arg1: verification (0=no, 1=yes)
\n
"
<<
"arg2: initialization (0=no init, 1=integer value, 2=decimal value)
\n
"
<<
"arg3: time kernel (0=no, 1=yes)
\n
"
<<
ck
::
utils
::
conv
::
get_conv_param_parser_helper_msg
()
<<
std
::
endl
;
}
bool
run_convnd_fwd_example
(
int
argc
,
char
*
argv
[])
{
print_helper_msg
();
...
...
@@ -79,5 +87,5 @@ bool run_convnd_fwd_example(int argc, char* argv[])
return
run
();
}
return
tru
e
;
return
fals
e
;
}
include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp
View file @
fc870c69
...
...
@@ -451,8 +451,8 @@ struct Sigmoid
is_same
<
T
,
ck
::
half_t
>::
value
||
is_same
<
T
,
int8_t
>::
value
||
is_same
<
T
,
int32_t
>::
value
,
"Data type is not supported by this operation!"
);
constexpr
T
one
{
1
}
;
y
=
one
/
(
ck
::
type_convert
<
T
>
(
one
)
+
ck
::
math
::
exp
(
-
x
));
constexpr
T
one
=
type_convert
<
T
>
(
1
)
;
y
=
one
/
(
ck
::
type_convert
<
T
>
(
one
)
+
ck
::
math
::
exp
(
-
x
));
};
};
...
...
@@ -489,7 +489,7 @@ struct Swish
y
=
type_convert
<
Y
>
(
x
/
(
1.
f
+
ck
::
math
::
exp
(
bx
)));
};
float
beta_
=
1.0
f
;
const
float
beta_
;
};
struct
SoftRelu
...
...
@@ -503,11 +503,11 @@ struct SoftRelu
is_same
<
T
,
half_t
>::
value
||
is_same
<
T
,
int32_t
>::
value
||
is_same
<
T
,
int8_t
>::
value
,
"Data type is not supported by this operation!"
);
T
casted_alpha
=
type_convert
<
T
>
(
alpha_
);
constexpr
T
one
{
1
}
;
y
=
ck
::
math
::
log
(
one
+
ck
::
math
::
exp
(
x
*
casted_alpha
))
/
casted_alpha
;
T
casted_alpha
=
type_convert
<
T
>
(
alpha_
);
constexpr
T
one
=
type_convert
<
T
>
(
1
)
;
y
=
ck
::
math
::
log
(
one
+
ck
::
math
::
exp
(
x
*
casted_alpha
))
/
casted_alpha
;
}
float
alpha_
;
const
float
alpha_
;
};
struct
Power
...
...
@@ -528,9 +528,9 @@ struct Power
T
shifted_scaled_x
=
casted_alpha
+
casted_beta
*
x
;
y
=
ck
::
math
::
pow
(
shifted_scaled_x
,
casted_gamma
);
}
float
alpha_
;
float
beta_
;
float
gamma_
;
const
float
alpha_
;
const
float
beta_
;
const
float
gamma_
;
};
struct
ClippedRelu
...
...
@@ -548,8 +548,8 @@ struct ClippedRelu
T
casted_beta
=
type_convert
<
T
>
(
beta_
);
y
=
ck
::
math
::
min
(
casted_beta
,
ck
::
math
::
max
(
casted_alpha
,
x
));
}
float
alpha_
;
float
beta_
;
const
float
alpha_
;
const
float
beta_
;
};
struct
LeakyRelu
...
...
@@ -566,7 +566,7 @@ struct LeakyRelu
T
casted_alpha
=
type_convert
<
T
>
(
alpha_
);
y
=
x
>=
0
?
x
:
x
*
casted_alpha
;
}
float
alpha_
;
const
float
alpha_
;
};
struct
Elu
...
...
@@ -583,7 +583,7 @@ struct Elu
T
casted_alpha
=
type_convert
<
T
>
(
alpha_
);
y
=
x
>
0
?
x
:
casted_alpha
*
ck
::
math
::
expm1
(
x
);
}
float
alpha_
;
const
float
alpha_
;
};
}
// namespace element_wise
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment