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
2f24b0eb
"git@developer.sourcefind.cn:modelzoo/resnet50_tensorflow.git" did not exist on "b94c9598889142463926b27edd81aa2afd67d7fe"
Commit
2f24b0eb
authored
Mar 24, 2023
by
Po-Yen, Chen
Browse files
Add instance for layout=NN
parent
552c6b23
Changes
3
Expand all
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
57 additions
and
43 deletions
+57
-43
library/include/ck/library/tensor_operation_instance/gpu/gemm.hpp
...include/ck/library/tensor_operation_instance/gpu/gemm.hpp
+1
-1
library/src/tensor_operation_instance/gpu/gemm/CMakeLists.txt
...ary/src/tensor_operation_instance/gpu/gemm/CMakeLists.txt
+5
-1
library/src/tensor_operation_instance/gpu/gemm/device_gemm_xdl_f16_f16_f16_km_nk_mn_instance.cpp
...pu/gemm/device_gemm_xdl_f16_f16_f16_km_nk_mn_instance.cpp
+51
-41
No files found.
library/include/ck/library/tensor_operation_instance/gpu/gemm.hpp
View file @
2f24b0eb
...
@@ -316,7 +316,7 @@ struct DeviceOperationInstanceFactory<
...
@@ -316,7 +316,7 @@ struct DeviceOperationInstanceFactory<
else
if
constexpr
(
is_same_v
<
ALayout
,
Col
>
&&
is_same_v
<
BLayout
,
Col
>
&&
else
if
constexpr
(
is_same_v
<
ALayout
,
Col
>
&&
is_same_v
<
BLayout
,
Col
>
&&
is_same_v
<
CLayout
,
Row
>
)
is_same_v
<
CLayout
,
Row
>
)
{
{
//
add_device_gemm_xdl_f16_f16_f16_km_nk_mn_instances(op_ptrs);
add_device_gemm_xdl_f16_f16_f16_km_nk_mn_instances
(
op_ptrs
);
// add_device_gemm_dl_f16_f16_f16_km_nk_mn_instances(op_ptrs);
// add_device_gemm_dl_f16_f16_f16_km_nk_mn_instances(op_ptrs);
// add_device_gemm_xdl_c_shuffle_f16_f16_f16_km_nk_mn_instances(op_ptrs);
// add_device_gemm_xdl_c_shuffle_f16_f16_f16_km_nk_mn_instances(op_ptrs);
}
}
...
...
library/src/tensor_operation_instance/gpu/gemm/CMakeLists.txt
View file @
2f24b0eb
...
@@ -10,7 +10,7 @@ add_instance_library(device_gemm_instance
...
@@ -10,7 +10,7 @@ add_instance_library(device_gemm_instance
device_gemm_xdl_f16_f16_f16_mk_kn_mn_instance.cpp
device_gemm_xdl_f16_f16_f16_mk_kn_mn_instance.cpp
device_gemm_xdl_f16_f16_f16_mk_nk_mn_instance.cpp
device_gemm_xdl_f16_f16_f16_mk_nk_mn_instance.cpp
# device_gemm_xdl_f16_f16_f16_km_kn_mn_instance.cpp
# device_gemm_xdl_f16_f16_f16_km_kn_mn_instance.cpp
#
device_gemm_xdl_f16_f16_f16_km_nk_mn_instance.cpp
device_gemm_xdl_f16_f16_f16_km_nk_mn_instance.cpp
# device_gemm_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instance.cpp
# device_gemm_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instance.cpp
# device_gemm_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instance.cpp
# device_gemm_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instance.cpp
# device_gemm_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instance.cpp
# device_gemm_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instance.cpp
...
@@ -44,6 +44,10 @@ add_instance_library(device_gemm_instance
...
@@ -44,6 +44,10 @@ add_instance_library(device_gemm_instance
set
(
ENABLE_IGLP_OPT ON
)
set
(
ENABLE_IGLP_OPT ON
)
if
(
ENABLE_IGLP_OPT
)
if
(
ENABLE_IGLP_OPT
)
# layout=NN
set_source_files_properties
(
device_gemm_xdl_f16_f16_f16_km_nk_mn_instance.cpp PROPERTIES
COMPILE_OPTIONS
";--save-temps;-Wno-gnu-line-marker;"
COMPILE_DEFINITIONS
";USE_IGLP_OPT;USE_WAVES_PER_EU;"
)
# layout=TT
# layout=TT
set_source_files_properties
(
device_gemm_xdl_f16_f16_f16_mk_kn_mn_instance.cpp PROPERTIES
set_source_files_properties
(
device_gemm_xdl_f16_f16_f16_mk_kn_mn_instance.cpp PROPERTIES
COMPILE_OPTIONS
";--save-temps;-Wno-gnu-line-marker;"
COMPILE_OPTIONS
";--save-temps;-Wno-gnu-line-marker;"
...
...
library/src/tensor_operation_instance/gpu/gemm/device_gemm_xdl_f16_f16_f16_km_nk_mn_instance.cpp
View file @
2f24b0eb
This diff is collapsed.
Click to expand it.
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