FazBrowse GitHub Viewer
|
Trending
|
URL:
|
Home
Tools:
[Download Repo ZIP]
[View Raw Code]
[Original HTTPS Page]
composable_kernel/example/01_gemm/gemm_dl_int4.cpp at develop · ROCm/composable_kernel · GitHub
Uh oh!
There was an error while loading.
Please reload this page
.
ROCm
/
composable_kernel
Public
Notifications
You must be signed in to change notification settings
Fork
312
Star
548
Code
Issues
9
Pull requests
8
Discussions
Actions
Projects
Wiki
Security and quality
0
Insights
Additional navigation options
Code
Issues
Pull requests
Discussions
Actions
Projects
Wiki
Security and quality
Insights
Expand file tree
Breadcrumbs
composable_kernel
/
example
/
01_gemm
/
gemm_dl_int4.cpp
Copy path
More file actions
More file actions
Latest commit
History
History
History
44 lines (33 loc) · 3.96 KB
Breadcrumbs
composable_kernel
/
example
/
01_gemm
/
gemm_dl_int4.cpp
Copy path
File metadata and controls
44 lines (33 loc) · 3.96 KB
Raw
Copy raw file
Download raw file
Open symbols panel
Edit and raw actions
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
//
Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
//
SPDX-License-Identifier: MIT
#
ifdef
CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
#
include
"
common.hpp
"
#
include
"
ck/tensor_operation/gpu/device/impl/device_gemm_dl.hpp
"
using
ADataType = ck::
int4_t
;
using
BDataType = ck::
int4_t
;
using
CDataType = ck::
int4_t
;
using
KernelADataType =
int8_t
;
using
KernelBDataType =
int8_t
;
using
KernelCDataType =
int8_t
;
using
AccDataType =
int32_t
;
using
ALayout = Col;
using
BLayout = Row;
using
CLayout = Row;
using
AElementOp = PassThrough;
using
BElementOp = PassThrough;
using
CElementOp = PassThrough;
static
constexpr
auto
GemmDefault = ck::tensor_operation::device::GemmSpecialization::Default;
//
clang-format off
using
DeviceGemmInstance = ck::tensor_operation::device::DeviceGemmDl
//
######| AData| BData| CData| AccData| ALayout| BLayout| CLayout| A| B| C| GEMM| Block| MPer| NPer| K0Per| K1| M1Per| N1Per| KPer| M11N11Thread| M11N11Thread| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| ABlockTransfer| BBlockTransfer| BBlockTransfer| BBlockTransfer| BBlockTransfer| BBlockTransfer| BBlockTransfer| BBlockTransfer| CThreadTransfer| CThreadTransfer| CThreadTransfer|
//
######| Type| Type| Type| Type| | | | Elementwise| Elementwise| Elementwise| Spacialization| Size| Block| Block| Block| | ThreadM111| ThreadN111| Thread| ClusterM110Xs| ClusterN110Xs| ThreadSliceLengths| ThreadClusterLengths| ThreadCluster| SrcAccess| SrcVectorTensor| SrcVectorTensor| DstVectorTensor| ThreadSliceLengths| ThreadClusterLengths| ThreadCluster| SrcAccess| SrcVectorTensor| SrcVectorTensor| DstVectorTensor| SrcDstAccess| SrcDstVectorDim| DstScalarPerVector|
//
######| | | | | | | | Operation| Operation| Operation| | | | | | | | | | | | K0_M0_M1_K1| K0_M0_M1_K1| ArrangeOrder| Order| Lengths_K0_M0_M1_K1| ContiguousDimOrder| Lengths_K0_M0_M1_K1| K0_N0_N1_K1| K0_N0_N1_K1| ArrangeOrder| Order| Lengths_K0_N0_N1_K1| ContiguousDimOrder| Lengths_K0_N0_N1_K1| Order| | |
//
######| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | |
< KernelADataType, KernelBDataType, KernelCDataType, AccDataType, ALayout, BLayout, CLayout, AElementOp, BElementOp, CElementOp, GemmDefault,
256
,
128
,
128
,
16
,
4
,
4
,
4
,
1
, S<
8
,
2
>, S<
8
,
2
>, S<
2
,
1
,
4
,
4
>, S<
8
,
1
,
32
,
1
>, S<
0
,
3
,
1
,
2
>, S<
0
,
3
,
1
,
2
>, S<
1
,
1
,
4
,
1
>, S<
0
,
3
,
1
,
2
>, S<
1
,
1
,
4
,
4
>, S<
2
,
1
,
4
,
4
>, S<
8
,
1
,
32
,
1
>, S<
0
,
3
,
1
,
2
>, S<
0
,
3
,
1
,
2
>, S<
1
,
1
,
4
,
1
>, S<
0
,
3
,
1
,
2
>, S<
1
,
1
,
4
,
4
>, S<
0
,
1
,
2
,
3
,
4
,
5
>,
5
,
4
>;
//
clang-format on
using
ReferenceGemmInstance = ck::tensor_operation::host::
ReferenceGemm<ADataType, BDataType, CDataType, AccDataType, AElementOp, BElementOp, CElementOp>;
#
define
BUILD_INT4_EXAMPLE
#
include
"
run_gemm_example.inc
"
int
main
(
int
argc,
char
* argv[]) {
return
!
run_gemm_example
(argc, argv); }
#
endif
Back
|
FazBrowse Home
|
New Git URL