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
5e8ee548
Commit
5e8ee548
authored
Mar 24, 2023
by
Po-Yen, Chen
Browse files
Add instance for layout=NT
parent
2f24b0eb
Changes
3
Expand all
Show whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
59 additions
and
45 deletions
+59
-45
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_kn_mn_instance.cpp
...pu/gemm/device_gemm_xdl_f16_f16_f16_km_kn_mn_instance.cpp
+53
-43
No files found.
library/include/ck/library/tensor_operation_instance/gpu/gemm.hpp
View file @
5e8ee548
...
@@ -309,7 +309,7 @@ struct DeviceOperationInstanceFactory<
...
@@ -309,7 +309,7 @@ struct DeviceOperationInstanceFactory<
else
if
constexpr
(
is_same_v
<
ALayout
,
Col
>
&&
is_same_v
<
BLayout
,
Row
>
&&
else
if
constexpr
(
is_same_v
<
ALayout
,
Col
>
&&
is_same_v
<
BLayout
,
Row
>
&&
is_same_v
<
CLayout
,
Row
>
)
is_same_v
<
CLayout
,
Row
>
)
{
{
//
add_device_gemm_xdl_f16_f16_f16_km_kn_mn_instances(op_ptrs);
add_device_gemm_xdl_f16_f16_f16_km_kn_mn_instances
(
op_ptrs
);
// add_device_gemm_dl_f16_f16_f16_km_kn_mn_instances(op_ptrs);
// add_device_gemm_dl_f16_f16_f16_km_kn_mn_instances(op_ptrs);
// add_device_gemm_xdl_c_shuffle_f16_f16_f16_km_kn_mn_instances(op_ptrs);
// add_device_gemm_xdl_c_shuffle_f16_f16_f16_km_kn_mn_instances(op_ptrs);
}
}
...
...
library/src/tensor_operation_instance/gpu/gemm/CMakeLists.txt
View file @
5e8ee548
...
@@ -9,7 +9,7 @@ add_instance_library(device_gemm_instance
...
@@ -9,7 +9,7 @@ add_instance_library(device_gemm_instance
# device_gemm_xdl_f32_f32_f32_km_nk_mn_instance.cpp
# device_gemm_xdl_f32_f32_f32_km_nk_mn_instance.cpp
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
...
@@ -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=NT
set_source_files_properties
(
device_gemm_xdl_f16_f16_f16_km_kn_mn_instance.cpp PROPERTIES
COMPILE_OPTIONS
";--save-temps;-Wno-gnu-line-marker;"
COMPILE_DEFINITIONS
";USE_IGLP_OPT;USE_WAVES_PER_EU;"
)
# layout=NN
# layout=NN
set_source_files_properties
(
device_gemm_xdl_f16_f16_f16_km_nk_mn_instance.cpp PROPERTIES
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_OPTIONS
";--save-temps;-Wno-gnu-line-marker;"
...
...
library/src/tensor_operation_instance/gpu/gemm/device_gemm_xdl_f16_f16_f16_km_kn_mn_instance.cpp
View file @
5e8ee548
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