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
629257f9
Commit
629257f9
authored
Oct 16, 2024
by
rocking
Browse files
Remove fp32 instances
parent
ba5d34aa
Changes
3
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
8 additions
and
1 deletion
+8
-1
example/ck_tile/02_layernorm2d/example_layernorm2d_fwd.cpp
example/ck_tile/02_layernorm2d/example_layernorm2d_fwd.cpp
+1
-1
example/ck_tile/02_layernorm2d/instances/layernorm2d_fwd_fp32_instance.cpp
...2_layernorm2d/instances/layernorm2d_fwd_fp32_instance.cpp
+2
-0
example/ck_tile/02_layernorm2d/layernorm2d_fwd_api.cpp
example/ck_tile/02_layernorm2d/layernorm2d_fwd_api.cpp
+5
-0
No files found.
example/ck_tile/02_layernorm2d/example_layernorm2d_fwd.cpp
View file @
629257f9
...
@@ -9,7 +9,7 @@ auto create_args(int argc, char* argv[])
...
@@ -9,7 +9,7 @@ auto create_args(int argc, char* argv[])
.
insert
(
"n"
,
"4096"
,
"m dimension"
)
.
insert
(
"n"
,
"4096"
,
"m dimension"
)
.
insert
(
"e"
,
"1e-5"
,
"epsilon"
)
.
insert
(
"e"
,
"1e-5"
,
"epsilon"
)
.
insert
(
"v"
,
"1"
,
"cpu validation or not"
)
.
insert
(
"v"
,
"1"
,
"cpu validation or not"
)
.
insert
(
"prec"
,
"fp
32
"
,
"precision"
)
.
insert
(
"prec"
,
"fp
16
"
,
"precision"
)
.
insert
(
"warmup"
,
"5"
,
"cold iter"
)
.
insert
(
"warmup"
,
"5"
,
"cold iter"
)
.
insert
(
"repeat"
,
"20"
,
"hot iter"
);
.
insert
(
"repeat"
,
"20"
,
"hot iter"
);
...
...
example/ck_tile/02_layernorm2d/instances/layernorm2d_fwd_fp32_instance.cpp
View file @
629257f9
...
@@ -7,6 +7,7 @@
...
@@ -7,6 +7,7 @@
#include "layernorm_dispatch.hpp"
#include "layernorm_dispatch.hpp"
// clang-format off
// clang-format off
#ifdef CK_TILE_LAYERNORM2D_FWD_FP32_DEFAULT
template
float
run_layernorm
<
float
,
1
,
32
,
4
,
false
>(
const
layernorm2d_fwd_args
&
param
,
ck_tile
::
stream_config
stream
);
template
float
run_layernorm
<
float
,
1
,
32
,
4
,
false
>(
const
layernorm2d_fwd_args
&
param
,
ck_tile
::
stream_config
stream
);
template
float
run_layernorm
<
float
,
1
,
64
,
2
,
false
>(
const
layernorm2d_fwd_args
&
param
,
ck_tile
::
stream_config
stream
);
template
float
run_layernorm
<
float
,
1
,
64
,
2
,
false
>(
const
layernorm2d_fwd_args
&
param
,
ck_tile
::
stream_config
stream
);
template
float
run_layernorm
<
float
,
1
,
64
,
4
,
false
>(
const
layernorm2d_fwd_args
&
param
,
ck_tile
::
stream_config
stream
);
template
float
run_layernorm
<
float
,
1
,
64
,
4
,
false
>(
const
layernorm2d_fwd_args
&
param
,
ck_tile
::
stream_config
stream
);
...
@@ -31,4 +32,5 @@ template float run_layernorm<float, 8, 64, 4, true>(const layernorm2d_fwd_args&
...
@@ -31,4 +32,5 @@ template float run_layernorm<float, 8, 64, 4, true>(const layernorm2d_fwd_args&
template
float
run_layernorm
<
float
,
16
,
64
,
2
,
true
>(
const
layernorm2d_fwd_args
&
param
,
ck_tile
::
stream_config
stream
);
template
float
run_layernorm
<
float
,
16
,
64
,
2
,
true
>(
const
layernorm2d_fwd_args
&
param
,
ck_tile
::
stream_config
stream
);
template
float
run_layernorm
<
float
,
8
,
64
,
4
,
true
,
true
>(
const
layernorm2d_fwd_args
&
param
,
ck_tile
::
stream_config
stream
);
template
float
run_layernorm
<
float
,
8
,
64
,
4
,
true
,
true
>(
const
layernorm2d_fwd_args
&
param
,
ck_tile
::
stream_config
stream
);
template
float
run_layernorm
<
float
,
16
,
64
,
2
,
true
,
true
>(
const
layernorm2d_fwd_args
&
param
,
ck_tile
::
stream_config
stream
);
template
float
run_layernorm
<
float
,
16
,
64
,
2
,
true
,
true
>(
const
layernorm2d_fwd_args
&
param
,
ck_tile
::
stream_config
stream
);
#endif
// clang-format on
// clang-format on
example/ck_tile/02_layernorm2d/layernorm2d_fwd_api.cpp
View file @
629257f9
...
@@ -112,6 +112,7 @@ float layernorm2d_fwd(layernorm2d_fwd_traits t,
...
@@ -112,6 +112,7 @@ float layernorm2d_fwd(layernorm2d_fwd_traits t,
}
}
}
}
}
}
#ifdef CK_TILE_LAYERNORM2D_FWD_FP32_DEFAULT
else
if
(
t
.
data_type
.
compare
(
"fp32"
)
==
0
)
else
if
(
t
.
data_type
.
compare
(
"fp32"
)
==
0
)
{
{
if
(
a
.
N
%
4
==
0
)
if
(
a
.
N
%
4
==
0
)
...
@@ -181,6 +182,10 @@ float layernorm2d_fwd(layernorm2d_fwd_traits t,
...
@@ -181,6 +182,10 @@ float layernorm2d_fwd(layernorm2d_fwd_traits t,
}
}
}
}
}
}
#endif
if
(
r
<
0
)
throw
std
::
runtime_error
(
"Without supported instances!"
);
return
r
;
return
r
;
}
}
...
...
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