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
da1495f7
"vscode:/vscode.git/clone" did not exist on "f911363143120049453e770ff850fc859387b6c2"
Commit
da1495f7
authored
Mar 15, 2023
by
rocking
Browse files
Separete different layout to different cpp
parent
ca7015c4
Changes
11
Show whitespace changes
Inline
Side-by-side
Showing
11 changed files
with
278 additions
and
154 deletions
+278
-154
library/src/tensor_operation_instance/gpu/quantization/CMakeLists.txt
...tensor_operation_instance/gpu/quantization/CMakeLists.txt
+8
-2
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_instance.hpp
...vice_gemm_quantization_dl_c_shuffle_i8_i8_i8_instance.hpp
+0
-68
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_km_kn_mn_instance.cpp
..._quantization_dl_c_shuffle_i8_i8_i8_km_kn_mn_instance.cpp
+32
-0
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_km_nk_mn_instance.cpp
..._quantization_dl_c_shuffle_i8_i8_i8_km_nk_mn_instance.cpp
+32
-0
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_kn_mn_instance.cpp
..._quantization_dl_c_shuffle_i8_i8_i8_mk_kn_mn_instance.cpp
+32
-0
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_nk_mn_instance.cpp
..._quantization_dl_c_shuffle_i8_i8_i8_mk_nk_mn_instance.cpp
+32
-0
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_instance.hpp
...ice_gemm_quantization_xdl_c_shuffle_i8_i8_i8_instance.hpp
+0
-84
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instance.cpp
...quantization_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instance.cpp
+35
-0
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_km_nk_mn_instance.cpp
...quantization_xdl_c_shuffle_i8_i8_i8_km_nk_mn_instance.cpp
+35
-0
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instance.cpp
...quantization_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instance.cpp
+36
-0
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instance.cpp
...quantization_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instance.cpp
+36
-0
No files found.
library/src/tensor_operation_instance/gpu/quantization/CMakeLists.txt
View file @
da1495f7
...
...
@@ -19,8 +19,14 @@ set(CONV2D_BIAS_PERCHANNEL_QUANT_SRC
)
set
(
GEMM_QUANT_SRC
gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_instance.cpp
gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_instance.cpp
gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_km_kn_mn_instance.cpp
gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_km_nk_mn_instance.cpp
gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_kn_mn_instance.cpp
gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_nk_mn_instance.cpp
gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instance.cpp
gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_km_nk_mn_instance.cpp
gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instance.cpp
gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instance.cpp
)
add_instance_library
(
device_quantization_instance
...
...
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_instance.
c
pp
→
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_instance.
h
pp
View file @
da1495f7
...
...
@@ -53,74 +53,6 @@ using device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_nk_mn_instances = std::t
// clang-format on
>
;
void
add_device_gemm_quantization_dl_c_shuffle_i8_i8_i8_km_kn_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Col
,
Row
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_dl_c_shuffle_i8_i8_i8_km_kn_mn_instances
<
Mul_Clamp
>
{});
}
void
add_device_gemm_quantization_dl_c_shuffle_i8_i8_i8_km_nk_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Col
,
Col
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_dl_c_shuffle_i8_i8_i8_km_nk_mn_instances
<
Mul_Clamp
>
{});
}
void
add_device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_kn_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Row
,
Row
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_kn_mn_instances
<
Mul_Clamp
>
{});
}
void
add_device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_nk_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Row
,
Col
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_nk_mn_instances
<
Mul_Clamp
>
{});
}
}
// namespace instance
}
// namespace device
}
// namespace tensor_operation
...
...
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_km_kn_mn_instance.cpp
0 → 100644
View file @
da1495f7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "device_gemm_quantization_dl_c_shuffle_i8_i8_i8_instance.hpp"
namespace
ck
{
namespace
tensor_operation
{
namespace
device
{
namespace
instance
{
// Layout(A, B, C) = [Col, Row, Row]
void
add_device_gemm_quantization_dl_c_shuffle_i8_i8_i8_km_kn_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Col
,
Row
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_dl_c_shuffle_i8_i8_i8_km_kn_mn_instances
<
Mul_Clamp
>
{});
}
}
// namespace instance
}
// namespace device
}
// namespace tensor_operation
}
// namespace ck
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_km_nk_mn_instance.cpp
0 → 100644
View file @
da1495f7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "device_gemm_quantization_dl_c_shuffle_i8_i8_i8_instance.hpp"
namespace
ck
{
namespace
tensor_operation
{
namespace
device
{
namespace
instance
{
// Layout(A, B, C) = [Col, Col, Row]
void
add_device_gemm_quantization_dl_c_shuffle_i8_i8_i8_km_nk_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Col
,
Col
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_dl_c_shuffle_i8_i8_i8_km_nk_mn_instances
<
Mul_Clamp
>
{});
}
}
// namespace instance
}
// namespace device
}
// namespace tensor_operation
}
// namespace ck
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_kn_mn_instance.cpp
0 → 100644
View file @
da1495f7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "device_gemm_quantization_dl_c_shuffle_i8_i8_i8_instance.hpp"
namespace
ck
{
namespace
tensor_operation
{
namespace
device
{
namespace
instance
{
// Layout(A, B, C) = [Row, Row, Row]
void
add_device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_kn_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Row
,
Row
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_kn_mn_instances
<
Mul_Clamp
>
{});
}
}
// namespace instance
}
// namespace device
}
// namespace tensor_operation
}
// namespace ck
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_nk_mn_instance.cpp
0 → 100644
View file @
da1495f7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "device_gemm_quantization_dl_c_shuffle_i8_i8_i8_instance.hpp"
namespace
ck
{
namespace
tensor_operation
{
namespace
device
{
namespace
instance
{
// Layout(A, B, C) = [Row, Col, Row]
void
add_device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_nk_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Row
,
Col
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_dl_c_shuffle_i8_i8_i8_mk_nk_mn_instances
<
Mul_Clamp
>
{});
}
}
// namespace instance
}
// namespace device
}
// namespace tensor_operation
}
// namespace ck
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_instance.
c
pp
→
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_instance.
h
pp
View file @
da1495f7
...
...
@@ -53,90 +53,6 @@ using device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instances = std::
// clang-format on
>
;
// Layout(A, B, C) = [Col, Row, Row]
void
add_device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Col
,
Row
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instances
<
Mul_Clamp
,
LoopScheduler
::
Default
,
PipelineVersion
::
v1
>
{});
}
// Layout(A, B, C) = [Col, Col, Row]
void
add_device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_km_nk_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Col
,
Col
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_km_nk_mn_instances
<
Mul_Clamp
,
LoopScheduler
::
Default
,
PipelineVersion
::
v1
>
{});
}
// Layout(A, B, C) = [Row, Row, Row]
void
add_device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Row
,
Row
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instances
<
Mul_Clamp
,
LoopScheduler
::
Default
,
PipelineVersion
::
v1
>
{});
}
// Layout(A, B, C) = [Row, Col, Row]
void
add_device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Row
,
Col
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instances
<
Mul_Clamp
,
LoopScheduler
::
Default
,
PipelineVersion
::
v1
>
{});
}
}
// namespace instance
}
// namespace device
}
// namespace tensor_operation
...
...
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instance.cpp
0 → 100644
View file @
da1495f7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_instance.hpp"
namespace
ck
{
namespace
tensor_operation
{
namespace
device
{
namespace
instance
{
// Layout(A, B, C) = [Col, Row, Row]
void
add_device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Col
,
Row
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instances
<
Mul_Clamp
,
LoopScheduler
::
Default
,
PipelineVersion
::
v1
>
{});
}
}
// namespace instance
}
// namespace device
}
// namespace tensor_operation
}
// namespace ck
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_km_nk_mn_instance.cpp
0 → 100644
View file @
da1495f7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_instance.hpp"
namespace
ck
{
namespace
tensor_operation
{
namespace
device
{
namespace
instance
{
// Layout(A, B, C) = [Col, Col, Row]
void
add_device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_km_nk_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Col
,
Col
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_km_nk_mn_instances
<
Mul_Clamp
,
LoopScheduler
::
Default
,
PipelineVersion
::
v1
>
{});
}
}
// namespace instance
}
// namespace device
}
// namespace tensor_operation
}
// namespace ck
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instance.cpp
0 → 100644
View file @
da1495f7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_instance.hpp"
namespace
ck
{
namespace
tensor_operation
{
namespace
device
{
namespace
instance
{
// Layout(A, B, C) = [Row, Row, Row]
void
add_device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Row
,
Row
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instances
<
Mul_Clamp
,
LoopScheduler
::
Default
,
PipelineVersion
::
v1
>
{});
}
}
// namespace instance
}
// namespace device
}
// namespace tensor_operation
}
// namespace ck
library/src/tensor_operation_instance/gpu/quantization/gemm/device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instance.cpp
0 → 100644
View file @
da1495f7
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_instance.hpp"
namespace
ck
{
namespace
tensor_operation
{
namespace
device
{
namespace
instance
{
// Layout(A, B, C) = [Row, Col, Row]
void
add_device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instances
(
std
::
vector
<
std
::
unique_ptr
<
DeviceGemmMultipleD
<
Row
,
Col
,
Empty_Tuple
,
Row
,
int8_t
,
int8_t
,
Empty_Tuple
,
int8_t
,
PassThrough
,
PassThrough
,
Mul_Clamp
>>>&
instances
)
{
add_device_operation_instances
(
instances
,
device_gemm_quantization_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instances
<
Mul_Clamp
,
LoopScheduler
::
Default
,
PipelineVersion
::
v1
>
{});
}
}
// namespace instance
}
// namespace device
}
// namespace tensor_operation
}
// namespace ck
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