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
62bc2102
Commit
62bc2102
authored
Dec 23, 2022
by
rocking
Browse files
Add comment
parent
07b9a7b1
Changes
1
Hide whitespace changes
Inline
Side-by-side
Showing
1 changed file
with
7 additions
and
2 deletions
+7
-2
include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_layernorm_xdl_cshuffle.hpp
...ce/impl/device_gemm_multiple_d_layernorm_xdl_cshuffle.hpp
+7
-2
No files found.
include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_layernorm_xdl_cshuffle.hpp
View file @
62bc2102
...
@@ -251,9 +251,14 @@ struct DeviceGemmMultipleDLayernorm_Xdl_CShuffle
...
@@ -251,9 +251,14 @@ struct DeviceGemmMultipleDLayernorm_Xdl_CShuffle
{
{
using
DeviceOp
=
DeviceGemmMultipleDLayernorm_Xdl_CShuffle
;
using
DeviceOp
=
DeviceGemmMultipleDLayernorm_Xdl_CShuffle
;
using
ELayout
=
HLayout
;
using
ELayout
=
HLayout
;
// EDataType, MeanDataType and VarDataType must be the same.
// eg. M, N, K = [1, 1, 1],
// in case of layernorm, divisor = 1 / sqrt(var + 1e-5) = 316.227783
// if (x - mean) != 0, (x - mean) * divisor * gamma might be too large
// However, (x - mean) * divisor * gamma should be 0 in this case
using
EDataType
=
HDataType
;
using
EDataType
=
HDataType
;
using
MeanDataType
=
CShuffle
DataType
;
using
MeanDataType
=
H
DataType
;
using
VarDataType
=
CShuffle
DataType
;
using
VarDataType
=
H
DataType
;
static
constexpr
index_t
NumDTensor
=
DsDataType
::
Size
();
static
constexpr
index_t
NumDTensor
=
DsDataType
::
Size
();
static
constexpr
index_t
LayernormHDstVectorSize
=
PostShuffleScalarPerVector
;
static
constexpr
index_t
LayernormHDstVectorSize
=
PostShuffleScalarPerVector
;
...
...
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