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_ROCM
Commits
1cb3e443
Commit
1cb3e443
authored
Oct 20, 2024
by
carlushuang
Browse files
rename more
parent
cc1898fc
Changes
6
Hide whitespace changes
Inline
Side-by-side
Showing
6 changed files
with
11 additions
and
11 deletions
+11
-11
example/ck_tile/02_layernorm2d/instances/layernorm2d_fwd_instance_common.hpp
...layernorm2d/instances/layernorm2d_fwd_instance_common.hpp
+2
-2
include/ck_tile/ops/layernorm2d.hpp
include/ck_tile/ops/layernorm2d.hpp
+3
-3
include/ck_tile/ops/layernorm2d/kernel/layernorm2d_fwd_shape.hpp
.../ck_tile/ops/layernorm2d/kernel/layernorm2d_fwd_shape.hpp
+1
-1
include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_rowwise_default_policy.hpp
...orm2d/pipeline/layernorm2d_fwd_rowwise_default_policy.hpp
+1
-1
include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_rowwise_pipeline.hpp
...layernorm2d/pipeline/layernorm2d_fwd_rowwise_pipeline.hpp
+3
-3
include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_rowwise_problem.hpp
.../layernorm2d/pipeline/layernorm2d_fwd_rowwise_problem.hpp
+1
-1
No files found.
example/ck_tile/02_layernorm2d/instances/layernorm2d_fwd_instance_common.hpp
View file @
1cb3e443
...
@@ -106,7 +106,7 @@ float layernorm2d_fwd_(const S& s, A a)
...
@@ -106,7 +106,7 @@ float layernorm2d_fwd_(const S& s, A a)
{
{
using
DataType
=
typename
Traits_
::
DataType
;
using
DataType
=
typename
Traits_
::
DataType
;
using
PipelineProblem
=
ck_tile
::
Layernorm2dFwd
WarpPerRow
Problem
<
using
PipelineProblem
=
ck_tile
::
Layernorm2dFwd
Rowwise
Problem
<
typename
LayerNormTypeConfig
<
DataType
>::
XDataType
,
typename
LayerNormTypeConfig
<
DataType
>::
XDataType
,
typename
LayerNormTypeConfig
<
DataType
>::
GammaDataType
,
typename
LayerNormTypeConfig
<
DataType
>::
GammaDataType
,
typename
LayerNormTypeConfig
<
DataType
>::
BetaDataType
,
typename
LayerNormTypeConfig
<
DataType
>::
BetaDataType
,
...
@@ -118,7 +118,7 @@ float layernorm2d_fwd_(const S& s, A a)
...
@@ -118,7 +118,7 @@ float layernorm2d_fwd_(const S& s, A a)
Traits_
::
kPadN
,
Traits_
::
kPadN
,
Traits_
::
kSaveMeanInvStd
,
Traits_
::
kSaveMeanInvStd
,
Traits_
::
kTwoPass
>
;
Traits_
::
kTwoPass
>
;
using
Pipeline
=
ck_tile
::
Layernorm2dFwd
WarpPerRow
Pipeline
<
PipelineProblem
>
;
using
Pipeline
=
ck_tile
::
Layernorm2dFwd
Rowwise
Pipeline
<
PipelineProblem
>
;
using
Kernel
=
ck_tile
::
Layernorm2dFwd
<
Pipeline
>
;
using
Kernel
=
ck_tile
::
Layernorm2dFwd
<
Pipeline
>
;
...
...
include/ck_tile/ops/layernorm2d.hpp
View file @
1cb3e443
...
@@ -5,7 +5,7 @@
...
@@ -5,7 +5,7 @@
#include "ck_tile/ops/layernorm2d/kernel/layernorm2d_fwd_kernel.hpp"
#include "ck_tile/ops/layernorm2d/kernel/layernorm2d_fwd_kernel.hpp"
#include "ck_tile/ops/layernorm2d/kernel/layernorm2d_fwd_shape.hpp"
#include "ck_tile/ops/layernorm2d/kernel/layernorm2d_fwd_shape.hpp"
#include "ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_
warp_per_row
_default_policy.hpp"
#include "ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_
rowwise
_default_policy.hpp"
#include "ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_
warp_per_row
_pipeline.hpp"
#include "ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_
rowwise
_pipeline.hpp"
#include "ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_
warp_per_row
_problem.hpp"
#include "ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_
rowwise
_problem.hpp"
#include "ck_tile/ops/common/tensor_layout.hpp"
#include "ck_tile/ops/common/tensor_layout.hpp"
include/ck_tile/ops/layernorm2d/kernel/layernorm2d_fwd_shape.hpp
View file @
1cb3e443
...
@@ -21,7 +21,7 @@ namespace ck_tile {
...
@@ -21,7 +21,7 @@ namespace ck_tile {
+--------------+--------------+ | <WarpPerBlock_M(2)> |
+--------------+--------------+ | <WarpPerBlock_M(2)> |
| wrap_2 | wrap_3 | | v
| wrap_2 | wrap_3 | | v
+--------------+--------------+--------------+--------------+----+ Block_M
+--------------+--------------+--------------+--------------+----+ Block_M
| | |
(Warp_M * WarpPerBlock_M * Repeat_M )
| | |
+ + |
+ + |
| | | v
| | | v
+--------------+--------------+--------------+--------------+ +
+--------------+--------------+--------------+--------------+ +
...
...
include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_
warp_per_row
_default_policy.hpp
→
include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_
rowwise
_default_policy.hpp
View file @
1cb3e443
...
@@ -9,7 +9,7 @@
...
@@ -9,7 +9,7 @@
namespace
ck_tile
{
namespace
ck_tile
{
struct
Layernorm2dFwd
WarpPerRow
DefaultPolicy
struct
Layernorm2dFwd
Rowwise
DefaultPolicy
{
{
template
<
typename
Problem
>
template
<
typename
Problem
>
CK_TILE_DEVICE
static
constexpr
auto
MakeXBlockTileDistribution
()
CK_TILE_DEVICE
static
constexpr
auto
MakeXBlockTileDistribution
()
...
...
include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_
warp_per_row
_pipeline.hpp
→
include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_
rowwise
_pipeline.hpp
View file @
1cb3e443
...
@@ -4,14 +4,14 @@
...
@@ -4,14 +4,14 @@
#pragma once
#pragma once
#include "ck_tile/core.hpp"
#include "ck_tile/core.hpp"
#include "ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_
warp_per_row
_default_policy.hpp"
#include "ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_
rowwise
_default_policy.hpp"
#include <string>
#include <string>
#include <type_traits>
#include <type_traits>
namespace
ck_tile
{
namespace
ck_tile
{
template
<
typename
Problem_
,
typename
Policy_
=
Layernorm2dFwd
WarpPerRow
DefaultPolicy
>
template
<
typename
Problem_
,
typename
Policy_
=
Layernorm2dFwd
Rowwise
DefaultPolicy
>
struct
Layernorm2dFwd
WarpPerRow
Pipeline
struct
Layernorm2dFwd
Rowwise
Pipeline
{
{
using
Problem
=
ck_tile
::
remove_cvref_t
<
Problem_
>
;
using
Problem
=
ck_tile
::
remove_cvref_t
<
Problem_
>
;
using
Policy
=
ck_tile
::
remove_cvref_t
<
Policy_
>
;
using
Policy
=
ck_tile
::
remove_cvref_t
<
Policy_
>
;
...
...
include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_
warp_per_row
_problem.hpp
→
include/ck_tile/ops/layernorm2d/pipeline/layernorm2d_fwd_
rowwise
_problem.hpp
View file @
1cb3e443
...
@@ -18,7 +18,7 @@ template <typename XDataType_,
...
@@ -18,7 +18,7 @@ template <typename XDataType_,
bool
kPadN_
,
bool
kPadN_
,
bool
kSaveMeanInvStd_
,
bool
kSaveMeanInvStd_
,
bool
kTwoPass_
>
bool
kTwoPass_
>
struct
Layernorm2dFwd
WarpPerRow
Problem
struct
Layernorm2dFwd
Rowwise
Problem
{
{
using
XDataType
=
remove_cvref_t
<
XDataType_
>
;
using
XDataType
=
remove_cvref_t
<
XDataType_
>
;
using
GammaDataType
=
remove_cvref_t
<
GammaDataType_
>
;
using
GammaDataType
=
remove_cvref_t
<
GammaDataType_
>
;
...
...
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