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
552c6b23
Commit
552c6b23
authored
Mar 24, 2023
by
Po-Yen, Chen
Browse files
Add instance for layout=TT
parent
82a7b60d
Changes
3
Expand all
Hide whitespace changes
Inline
Side-by-side
Showing
3 changed files
with
67 additions
and
62 deletions
+67
-62
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
+6
-1
library/src/tensor_operation_instance/gpu/gemm/device_gemm_xdl_f16_f16_f16_mk_kn_mn_instance.cpp
...pu/gemm/device_gemm_xdl_f16_f16_f16_mk_kn_mn_instance.cpp
+60
-60
No files found.
library/include/ck/library/tensor_operation_instance/gpu/gemm.hpp
View file @
552c6b23
...
...
@@ -295,7 +295,7 @@ struct DeviceOperationInstanceFactory<
if
constexpr
(
is_same_v
<
ALayout
,
Row
>
&&
is_same_v
<
BLayout
,
Row
>
&&
is_same_v
<
CLayout
,
Row
>
)
{
//
add_device_gemm_xdl_f16_f16_f16_mk_kn_mn_instances(op_ptrs);
add_device_gemm_xdl_f16_f16_f16_mk_kn_mn_instances
(
op_ptrs
);
// add_device_gemm_dl_f16_f16_f16_mk_kn_mn_instances(op_ptrs);
// add_device_gemm_xdl_c_shuffle_f16_f16_f16_mk_kn_mn_instances(op_ptrs);
}
...
...
library/src/tensor_operation_instance/gpu/gemm/CMakeLists.txt
View file @
552c6b23
...
...
@@ -7,7 +7,7 @@ add_instance_library(device_gemm_instance
# device_gemm_xdl_f32_f32_f32_mk_nk_mn_instance.cpp
# device_gemm_xdl_f32_f32_f32_km_kn_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_km_kn_mn_instance.cpp
# device_gemm_xdl_f16_f16_f16_km_nk_mn_instance.cpp
...
...
@@ -44,6 +44,11 @@ add_instance_library(device_gemm_instance
set
(
ENABLE_IGLP_OPT ON
)
if
(
ENABLE_IGLP_OPT
)
# layout=TT
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_DEFINITIONS
";USE_WAVES_PER_EU;"
)
# layout=TN
set_source_files_properties
(
device_gemm_xdl_f16_f16_f16_mk_nk_mn_instance.cpp PROPERTIES
COMPILE_OPTIONS
";--save-temps;-Wno-gnu-line-marker;-mllvm;-amdgpu-enable-max-ilp-scheduling-strategy;"
COMPILE_DEFINITIONS
";USE_IGLP_OPT;USE_WAVES_PER_EU;"
)
...
...
library/src/tensor_operation_instance/gpu/gemm/device_gemm_xdl_f16_f16_f16_mk_kn_mn_instance.cpp
View file @
552c6b23
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