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
128f0ee1
Commit
128f0ee1
authored
May 04, 2023
by
rocking
Browse files
Refine naming
parent
49eb83aa
Changes
3
Show whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
6 additions
and
6 deletions
+6
-6
example/27_layernorm/layernorm_fp16.cpp
example/27_layernorm/layernorm_fp16.cpp
+1
-1
example/27_layernorm/layernorm_splitk_fp16.cpp
example/27_layernorm/layernorm_splitk_fp16.cpp
+1
-1
include/ck/tensor_operation/gpu/device/impl/device_normalization_splitk_impl.hpp
...tion/gpu/device/impl/device_normalization_splitk_impl.hpp
+4
-4
No files found.
example/27_layernorm/layernorm_fp16.cpp
View file @
128f0ee1
...
@@ -27,7 +27,7 @@ using DeviceInstance =
...
@@ -27,7 +27,7 @@ using DeviceInstance =
32
,
// ClusterK
32
,
// ClusterK
1
,
// SliceM
1
,
// SliceM
8
,
// SliceK
8
,
// SliceK
1
,
//
SrcVec
Dim (0=M, 1=K)
1
,
//
XYVector
Dim (0=M, 1=K)
8
,
// SrcScalarPerVector
8
,
// SrcScalarPerVector
1
,
// GammaVecDim (0=M, 1=K)
1
,
// GammaVecDim (0=M, 1=K)
8
,
// GammaScalarPerVector
8
,
// GammaScalarPerVector
...
...
example/27_layernorm/layernorm_splitk_fp16.cpp
View file @
128f0ee1
...
@@ -27,7 +27,7 @@ using DeviceInstance =
...
@@ -27,7 +27,7 @@ using DeviceInstance =
32
,
// ClusterK
32
,
// ClusterK
1
,
// SliceM
1
,
// SliceM
8
,
// SliceK
8
,
// SliceK
1
,
//
SrcVec
Dim (0=M, 1=K)
1
,
//
XYVector
Dim (0=M, 1=K)
8
,
// XScalarPerVector
8
,
// XScalarPerVector
1
,
// GammaVecDim (0=M, 1=K)
1
,
// GammaVecDim (0=M, 1=K)
8
,
// GammaScalarPerVector
8
,
// GammaScalarPerVector
...
...
include/ck/tensor_operation/gpu/device/impl/device_normalization_splitk_impl.hpp
View file @
128f0ee1
...
@@ -221,17 +221,17 @@ struct DeviceNormalizationSplitKImpl : public DeviceNormalization<XDataType,
...
@@ -221,17 +221,17 @@ struct DeviceNormalizationSplitKImpl : public DeviceNormalization<XDataType,
template
<
typename
DoPads
,
index_t
MPerTile
,
index_t
KPerTile
>
template
<
typename
DoPads
,
index_t
MPerTile
,
index_t
KPerTile
>
static
auto
MakeMeanVarDescriptor_M_K
(
index_t
M
,
index_t
K
)
static
auto
MakeMeanVarDescriptor_M_K
(
index_t
M
,
index_t
K
)
{
{
const
auto
grid_desc_m_
n
=
const
auto
grid_desc_m_
k
=
make_naive_tensor_descriptor
(
make_tuple
(
M
,
K
),
make_tuple
(
K
,
I1
));
make_naive_tensor_descriptor
(
make_tuple
(
M
,
K
),
make_tuple
(
K
,
I1
));
return
PadTensorDescriptor
(
grid_desc_m_
n
,
make_tuple
(
MPerTile
,
KPerTile
),
DoPads
{});
return
PadTensorDescriptor
(
grid_desc_m_
k
,
make_tuple
(
MPerTile
,
KPerTile
),
DoPads
{});
}
}
template
<
typename
DoPads
,
index_t
MPerTile
,
index_t
KPerTile
>
template
<
typename
DoPads
,
index_t
MPerTile
,
index_t
KPerTile
>
static
auto
MakeCountDescriptor_M_K
(
index_t
M
,
index_t
K
)
static
auto
MakeCountDescriptor_M_K
(
index_t
M
,
index_t
K
)
{
{
const
auto
grid_desc_m_
n
=
const
auto
grid_desc_m_
k
=
make_naive_tensor_descriptor
(
make_tuple
(
M
,
K
),
make_tuple
(
I0
,
I1
));
make_naive_tensor_descriptor
(
make_tuple
(
M
,
K
),
make_tuple
(
I0
,
I1
));
return
PadTensorDescriptor
(
grid_desc_m_
n
,
make_tuple
(
MPerTile
,
KPerTile
),
DoPads
{});
return
PadTensorDescriptor
(
grid_desc_m_
k
,
make_tuple
(
MPerTile
,
KPerTile
),
DoPads
{});
}
}
using
SrcGridDesc_M_K
=
decltype
(
MakeSrc2dDescriptor
({
1
},
{
1
},
1
,
1
));
using
SrcGridDesc_M_K
=
decltype
(
MakeSrc2dDescriptor
({
1
},
{
1
},
1
,
1
));
...
...
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