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
95ad9091
Commit
95ad9091
authored
Jul 10, 2022
by
Wenkai
Browse files
Merge branch 'develop' into static_ck_small_gemm
parents
d2640676
1677cf70
Changes
44
Show whitespace changes
Inline
Side-by-side
Showing
20 changed files
with
497 additions
and
85 deletions
+497
-85
Jenkinsfile
Jenkinsfile
+44
-24
example/07_conv2d_fwd_bias_relu_add/conv2d_fwd_xdl_bias_relu_add.cpp
...conv2d_fwd_bias_relu_add/conv2d_fwd_xdl_bias_relu_add.cpp
+4
-4
example/12_reduce/reduce_blockwise.cpp
example/12_reduce/reduce_blockwise.cpp
+3
-3
example/12_reduce/reduce_blockwise_two_call.cpp
example/12_reduce/reduce_blockwise_two_call.cpp
+3
-3
example/13_pool2d_fwd/pool2d_fwd_common.hpp
example/13_pool2d_fwd/pool2d_fwd_common.hpp
+26
-20
example/13_pool2d_fwd/pool2d_fwd_fp16.cpp
example/13_pool2d_fwd/pool2d_fwd_fp16.cpp
+0
-2
example/13_pool2d_fwd/pool2d_fwd_fp32.cpp
example/13_pool2d_fwd/pool2d_fwd_fp32.cpp
+0
-2
example/15_grouped_gemm/grouped_gemm_xdl_fp16.cpp
example/15_grouped_gemm/grouped_gemm_xdl_fp16.cpp
+7
-2
example/16_gemm_reduce/gemm_reduce_xdl_max_fp16.cpp
example/16_gemm_reduce/gemm_reduce_xdl_max_fp16.cpp
+1
-1
example/16_gemm_reduce/gemm_reduce_xdl_mean_squaremean_fp16.cpp
...e/16_gemm_reduce/gemm_reduce_xdl_mean_squaremean_fp16.cpp
+2
-2
example/18_batched_gemm_reduce/batched_gemm_reduce_xdl_fp16.cpp
...e/18_batched_gemm_reduce/batched_gemm_reduce_xdl_fp16.cpp
+2
-2
example/19_binary_elementwise/broadcast_add_2d_amn_bn.cpp
example/19_binary_elementwise/broadcast_add_2d_amn_bn.cpp
+31
-5
example/19_binary_elementwise/broadcast_add_3d_am_bmnk.cpp
example/19_binary_elementwise/broadcast_add_3d_am_bmnk.cpp
+5
-4
example/19_binary_elementwise/elementwise_add_1d.cpp
example/19_binary_elementwise/elementwise_add_1d.cpp
+30
-4
example/19_binary_elementwise/elementwise_add_4d.cpp
example/19_binary_elementwise/elementwise_add_4d.cpp
+30
-4
example/21_gemm_layernorm/gemm_layernorm_xdl_fp16.cpp
example/21_gemm_layernorm/gemm_layernorm_xdl_fp16.cpp
+2
-2
example/22_cgemm/CMakeLists.txt
example/22_cgemm/CMakeLists.txt
+1
-0
example/22_cgemm/cgemm_xdl_fp16.cpp
example/22_cgemm/cgemm_xdl_fp16.cpp
+302
-0
example/CMakeLists.txt
example/CMakeLists.txt
+2
-1
include/ck/tensor_operation/gpu/device/device_base.hpp
include/ck/tensor_operation/gpu/device/device_base.hpp
+2
-0
No files found.
Jenkinsfile
View file @
95ad9091
...
...
@@ -212,30 +212,50 @@ def runCKProfiler(Map conf=[:]){
{
cmake_build
(
conf
)
dir
(
"script"
){
def
perf_log
=
"perf_gemm_${gpu_arch}.log"
sh
"rm -f ${perf_log}"
sh
"echo Branch name: ${env.BRANCH_NAME} > ${perf_log}"
sh
"./profile_gemm.sh gemm 0 0 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 1 0 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 2 0 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 3 0 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 0 1 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 1 1 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 2 1 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 3 1 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 0 2 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 1 2 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 2 2 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 3 2 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 0 3 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 1 3 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 2 3 0 1 0 5 | tee -a ${perf_log}"
sh
"./profile_gemm.sh gemm 3 3 0 1 0 5 | tee -a ${perf_log}"
//run gemm performance tests
def
gemm_log
=
"perf_gemm_${gpu_arch}.log"
sh
"rm -f ${gemm_log}"
sh
"echo Branch name: ${env.BRANCH_NAME} > ${gemm_log}"
sh
"echo Node name: ${NODE_NAME} >> ${gemm_log}"
sh
"echo GPU_arch: ${gpu_arch} >> ${gemm_log}"
sh
"hipcc --version | grep -e 'HIP version' >> ${gemm_log}"
sh
"/opt/rocm/bin/amdclang++ --version | grep -e 'InstalledDir' >> ${gemm_log}"
sh
"./profile_gemm.sh gemm 0 0 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 1 0 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 2 0 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 3 0 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 0 1 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 1 1 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 2 1 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 3 1 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 0 2 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 1 2 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 2 2 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 3 2 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 0 3 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 1 3 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 2 3 0 1 0 5 | tee -a ${gemm_log}"
sh
"./profile_gemm.sh gemm 3 3 0 1 0 5 | tee -a ${gemm_log}"
//results will be parsed, stored, and analyzed within the python script
//the script will return 0 if the performance criteria are met
//or return 1 if the criteria are not met
archiveArtifacts
"${perf_log}"
sh
"python3 parse_perf_data.py ${perf_log} "
archiveArtifacts
"${gemm_log}"
sh
"python3 parse_perf_data.py ${gemm_log} "
//run resnet50 test
def
resnet_log
=
"perf_resnet50_${gpu_arch}.log"
sh
"rm -f ${resnet_log}"
sh
"echo Branch name: ${env.BRANCH_NAME} > ${resnet_log}"
sh
"echo Node name: ${NODE_NAME} >> ${resnet_log}"
sh
"echo GPU_arch: ${gpu_arch} >> ${resnet_log}"
sh
"hipcc --version | grep -e 'HIP version' >> ${resnet_log}"
sh
"/opt/rocm/bin/amdclang++ --version | grep -e 'InstalledDir' >> ${resnet_log}"
//first run tests with N=256
sh
"./profile_conv.sh conv_fwd_bias_relu 1 1 1 1 0 2 0 1 256 | tee -a ${resnet_log}"
//then run with N=4
sh
"./profile_conv.sh conv_fwd_bias_relu 1 1 1 1 0 2 0 1 4 | tee -a ${resnet_log}"
archiveArtifacts
"${resnet_log}"
//the script will put the results from N=256 and N=4 runs into separate tables
sh
"python3 parse_perf_data.py ${resnet_log} "
}
}
}
...
...
example/07_conv2d_fwd_bias_relu_add/conv2d_fwd_xdl_bias_relu_add.cpp
View file @
95ad9091
...
...
@@ -224,10 +224,10 @@ int main(int argc, char* argv[])
{
case
0
:
break
;
case
1
:
input
.
GenerateTensorValue
(
GeneratorTensor_2
<
InDataType
>
{
-
5
,
5
});
weights
.
GenerateTensorValue
(
GeneratorTensor_2
<
WeiDataType
>
{
-
5
,
5
});
bias
.
GenerateTensorValue
(
GeneratorTensor_2
<
OutDataType
>
{
-
5
,
5
});
residual
.
GenerateTensorValue
(
GeneratorTensor_2
<
OutDataType
>
{
-
5
,
5
});
input
.
GenerateTensorValue
(
GeneratorTensor_2
<
InDataType
>
{
-
2
,
2
});
weights
.
GenerateTensorValue
(
GeneratorTensor_2
<
WeiDataType
>
{
-
2
,
2
});
bias
.
GenerateTensorValue
(
GeneratorTensor_2
<
OutDataType
>
{
-
2
,
2
});
residual
.
GenerateTensorValue
(
GeneratorTensor_2
<
OutDataType
>
{
-
2
,
2
});
break
;
default:
input
.
GenerateTensorValue
(
GeneratorTensor_3
<
InDataType
>
{
0.0
,
1.0
});
...
...
example/12_reduce/reduce_blockwise.cpp
View file @
95ad9091
...
...
@@ -147,8 +147,6 @@ class SimpleAppArgs
int
main
(
int
argc
,
char
*
argv
[])
{
using
namespace
ck
::
host_reduce
;
const
std
::
vector
<
int
>
reduceDims
{
0
,
1
,
2
};
const
std
::
vector
<
int
>
invariantDims
{
3
};
...
...
@@ -254,7 +252,9 @@ int main(int argc, char* argv[])
ReductionHost
<
InDataType
,
AccDataType
,
OutDataType
,
ReduceOpId
,
ReduceOperation
,
InElementwiseOperation
,
AccElementwiseOperation
,
Rank
,
NumReduceDim
,
PropagateNan
,
...
...
example/12_reduce/reduce_blockwise_two_call.cpp
View file @
95ad9091
...
...
@@ -108,8 +108,6 @@ int main(int argc, char* argv[])
const
std
::
vector
<
size_t
>
outLengths
=
{
64
,
320
,
80
};
using
namespace
ck
::
host_reduce
;
if
(
argc
==
1
)
{
do_verify
=
true
;
...
...
@@ -191,7 +189,9 @@ int main(int argc, char* argv[])
ReductionHost
<
InOutDataType
,
AccDataType
,
InOutDataType
,
ReduceOpId
,
ReduceOperation
,
InElementwiseOperation
,
AccElementwiseOperation
,
5
,
// Rank
2
,
// NumReduceDim
PropagateNan
,
...
...
example/13_pool2d_fwd/pool2d_fwd_common.hpp
View file @
95ad9091
...
...
@@ -8,10 +8,12 @@
#include "device.hpp"
#include "host_tensor.hpp"
#include "host_tensor_generator.hpp"
#include "host_reduce_util.hpp"
#include "device_tensor.hpp"
#include "tensor_layout.hpp"
#include "reduction_enums.hpp"
#include "reduction_operator_mapping.hpp"
#include "reduction_functions_accumulate.hpp"
#include "device_pool2d_fwd_nhwc_nhwc.hpp"
template
<
typename
InDataType
,
...
...
@@ -29,19 +31,24 @@ static void pool_host_verify(const Tensor<InDataType>& in,
const
std
::
array
<
ck
::
index_t
,
2
>&
in_left_pads
,
const
std
::
array
<
ck
::
index_t
,
2
>&
/*in_right_pads*/
)
{
using
namespace
ck
::
host_reduce
;
const
int32_t
divider
=
window_spatial_lengths
[
0
]
*
window_spatial_lengths
[
1
];
const
auto
PreUnaryOp
=
PreUnaryOpFn
<
AccDataType
,
ReduceOpId
>
(
divider
);
const
auto
PosUnaryOp
=
PosUnaryOpFn
<
AccDataType
,
ReduceOpId
>
(
divider
);
using
ReduceOperation
=
typename
ck
::
reduce_binary_operator
<
AccDataType
,
ReduceOpId
>::
opType
;
using
InElementwiseOperation
=
typename
ck
::
reduce_unary_operator
<
AccDataType
,
ReduceOpId
,
true
,
true
>::
InElementwiseOperation
;
using
AccElementwiseOperation
=
typename
ck
::
reduce_unary_operator
<
AccDataType
,
ReduceOpId
,
true
,
true
>::
AccElementwiseOperation
;
const
InElementwiseOperation
in_elementwise_op
(
divider
);
const
AccElementwiseOperation
acc_elementwise_op
(
divider
);
if
constexpr
(
!
OutputIndex
)
{
auto
opReduce
=
ReduceOpFn
<
AccDataType
,
ReduceOpId
>
();
using
Accumulation
=
ck
::
detail
::
AccumulateWithNanCheck
<
PropagateNan
,
ReduceOperation
,
AccDataType
>
;
auto
f_nchw
=
[
&
](
auto
n
,
auto
c
,
auto
ho
,
auto
wo
)
{
auto
accuVal
=
ReduceOp
Z
er
oVal
<
AccDataType
,
ReduceOpId
>
();
auto
accuVal
=
ReduceOper
ation
::
GetIdentityValue
();
for
(
ck
::
index_t
y
=
0
;
y
<
window_spatial_lengths
[
0
];
++
y
)
{
...
...
@@ -54,14 +61,14 @@ static void pool_host_verify(const Tensor<InDataType>& in,
{
AccDataType
currVal
=
static_cast
<
AccDataType
>
(
in
(
n
,
c
,
hi
,
wi
));
PreUnaryOp
(
currVal
);
in_elementwise_op
(
currVal
,
currVal
);
binop_with_nan_check
<
AccDataType
,
PropagateNan
>
(
opReduce
,
accuVal
,
currVal
);
Accumulation
::
Calculate
(
accuVal
,
currVal
);
}
}
}
PosUnaryOp
(
accuVal
);
acc_elementwise_op
(
accuVal
,
accuVal
);
out
(
n
,
c
,
ho
,
wo
)
=
accuVal
;
};
...
...
@@ -74,10 +81,12 @@ static void pool_host_verify(const Tensor<InDataType>& in,
}
else
{
auto
opReduce
=
ReduceOpFn2
<
AccDataType
,
ReduceOpId
>
();
using
Accumulation
=
ck
::
detail
::
AccumulateWithIndexAndNanCheck
<
PropagateNan
,
ReduceOperation
,
AccDataType
,
IndexDataType
>
;
auto
f_nchw
=
[
&
](
auto
n
,
auto
c
,
auto
ho
,
auto
wo
)
{
auto
accuVal
=
ReduceOp
Z
er
oVal
<
AccDataType
,
ReduceOpId
>
();
auto
accuVal
=
ReduceOper
ation
::
GetIdentityValue
();
IndexDataType
accuIndex
=
0
;
for
(
ck
::
index_t
y
=
0
;
y
<
window_spatial_lengths
[
0
];
++
y
)
...
...
@@ -92,15 +101,14 @@ static void pool_host_verify(const Tensor<InDataType>& in,
AccDataType
currVal
=
static_cast
<
AccDataType
>
(
in
(
n
,
c
,
hi
,
wi
));
IndexDataType
currIndex
=
y
*
window_spatial_lengths
[
1
]
+
x
;
PreUnaryOp
(
currVal
);
in_elementwise_op
(
currVal
,
currVal
);
binop_with_index_and_nan_check
<
AccDataType
,
IndexDataType
,
PropagateNan
>
(
opReduce
,
accuVal
,
currVal
,
accuIndex
,
currIndex
);
Accumulation
::
Calculate
(
accuVal
,
currVal
,
accuIndex
,
currIndex
);
}
}
}
PosUnaryOp
(
accuVal
);
acc_elementwise_op
(
accuVal
,
accuVal
);
out
(
n
,
c
,
ho
,
wo
)
=
accuVal
;
out_indices
(
n
,
c
,
ho
,
wo
)
=
accuIndex
;
...
...
@@ -139,8 +147,6 @@ bool pool_test(bool do_verification,
ck
::
index_t
in_right_pad_h
,
ck
::
index_t
in_right_pad_w
)
{
using
namespace
ck
::
host_reduce
;
using
DevicePoolFwdInstance
=
ck
::
tensor_operation
::
device
::
DevicePool2dFwd_Input_N_Hi_Wi_C_Output_N_Ho_Wo_C
<
InDataType
,
// InDataType
...
...
example/13_pool2d_fwd/pool2d_fwd_fp16.cpp
View file @
95ad9091
...
...
@@ -27,8 +27,6 @@ static constexpr bool PropagateNan = false;
int
main
(
int
argc
,
char
*
argv
[])
{
using
namespace
ck
::
host_reduce
;
bool
do_verification
;
int
init_method
;
bool
time_kernel
;
...
...
example/13_pool2d_fwd/pool2d_fwd_fp32.cpp
View file @
95ad9091
...
...
@@ -27,8 +27,6 @@ static constexpr bool PropagateNan = false;
int
main
(
int
argc
,
char
*
argv
[])
{
using
namespace
ck
::
host_reduce
;
bool
do_verification
;
int
init_method
;
bool
time_kernel
;
...
...
example/15_grouped_gemm/grouped_gemm_xdl_fp16.cpp
View file @
95ad9091
...
...
@@ -78,7 +78,7 @@ int main(int argc, char* argv[])
exit
(
0
);
}
int
group_count
=
4
;
int
group_count
=
rand
()
%
16
+
1
;
// GEMM shape
std
::
vector
<
ck
::
tensor_operation
::
device
::
GemmShape
>
gemm_shapes
;
...
...
@@ -189,12 +189,17 @@ int main(int argc, char* argv[])
auto
b_element_op
=
BElementOp
{};
auto
c_element_op
=
CElementOp
{};
// do GEMM
auto
gemm
=
DeviceGemmInstance
{};
auto
invoker
=
gemm
.
MakeInvoker
();
// do GEMM
auto
argument
=
gemm
.
MakeArgument
(
p_a
,
p_b
,
p_c
,
gemm_shapes
,
a_element_op
,
b_element_op
,
c_element_op
);
DeviceMem
gemm_desc_workspace
(
gemm
.
GetWorkSpaceSize
(
&
argument
));
gemm
.
SetWorkSpacePointer
(
&
argument
,
gemm_desc_workspace
.
GetDeviceBuffer
());
if
(
!
gemm
.
IsSupportedArgument
(
argument
))
{
throw
std
::
runtime_error
(
...
...
example/16_gemm_reduce/gemm_reduce_xdl_max_fp16.cpp
View file @
95ad9091
...
...
@@ -236,7 +236,7 @@ int main(int argc, char* argv[])
for
(
int
m
=
0
;
m
<
M
;
++
m
)
{
ReduceAccDataType
d_acc
=
d_reduce_op
.
Get
ReductionZero
Val
();
ReduceAccDataType
d_acc
=
d_reduce_op
.
Get
Identity
Val
ue
();
for
(
int
n
=
0
;
n
<
N
;
++
n
)
d_reduce_op
(
d_acc
,
c_m_n_host_result
(
m
,
n
));
...
...
example/16_gemm_reduce/gemm_reduce_xdl_mean_squaremean_fp16.cpp
View file @
95ad9091
...
...
@@ -261,8 +261,8 @@ int main(int argc, char* argv[])
for
(
int
m
=
0
;
m
<
M
;
++
m
)
{
float
d0_acc
=
d0_reduce_op
.
Get
ReductionZero
Val
();
float
d1_acc
=
d1_reduce_op
.
Get
ReductionZero
Val
();
float
d0_acc
=
d0_reduce_op
.
Get
Identity
Val
ue
();
float
d1_acc
=
d1_reduce_op
.
Get
Identity
Val
ue
();
for
(
int
n
=
0
;
n
<
N
;
++
n
)
{
...
...
example/18_batched_gemm_reduce/batched_gemm_reduce_xdl_fp16.cpp
View file @
95ad9091
...
...
@@ -259,8 +259,8 @@ int main(int argc, char* argv[])
{
for
(
int
m
=
0
;
m
<
M
;
++
m
)
{
float
d0_acc
=
d0_reduce_op
.
Get
ReductionZero
Val
();
float
d1_acc
=
d1_reduce_op
.
Get
ReductionZero
Val
();
float
d0_acc
=
d0_reduce_op
.
Get
Identity
Val
ue
();
float
d1_acc
=
d1_reduce_op
.
Get
Identity
Val
ue
();
for
(
int
n
=
0
;
n
<
N
;
++
n
)
{
...
...
example/19_binary_elementwise/broadcast_add_2d_amn_bn.cpp
View file @
95ad9091
/*******************************************************************************
*
* MIT License
*
* Copyright (c) 2022 Advanced Micro Devices, Inc.
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
* in the Software without restriction, including without limitation the rights
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
* copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in all
* copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*
*******************************************************************************/
#include <iostream>
#include <cstdlib>
#include "check_err.hpp"
...
...
@@ -17,7 +42,8 @@ using ABDataType = F16;
using
CDataType
=
F16
;
using
EltwiseComputeDataType
=
F32
;
using
Add
=
ck
::
tensor_operation
::
binary_element_wise
::
Add
;
using
Add
=
ck
::
tensor_operation
::
binary_element_wise
::
Add
<
EltwiseComputeDataType
,
EltwiseComputeDataType
,
EltwiseComputeDataType
>
;
using
DeviceElementwiseAddInstance
=
ck
::
tensor_operation
::
device
::
DeviceBinaryElementwise
<
ABDataType
,
...
...
@@ -46,19 +72,19 @@ void host_broadcast2D(
{
for
(
int
n
=
0
;
n
<
N
;
++
n
)
{
ComputeDataType
Amn
=
static_cas
t
<
ComputeDataType
>
(
A
(
m
,
n
));
ComputeDataType
Amn
=
ck
::
type_conver
t
<
ComputeDataType
>
(
A
(
m
,
n
));
ComputeDataType
Cmn
=
0
;
if
constexpr
(
broadcastDim
==
0
)
{
ComputeDataType
Bn
=
static_cas
t
<
ComputeDataType
>
(
B
(
n
));
ComputeDataType
Bn
=
ck
::
type_conver
t
<
ComputeDataType
>
(
B
(
n
));
functor
(
Cmn
,
Amn
,
Bn
);
}
else
{
ComputeDataType
Bm
=
static_cas
t
<
ComputeDataType
>
(
B
(
m
));
ComputeDataType
Bm
=
ck
::
type_conver
t
<
ComputeDataType
>
(
B
(
m
));
functor
(
Cmn
,
Amn
,
Bm
);
}
C
(
m
,
n
)
=
static_cas
t
<
ctype
>
(
Cmn
);
C
(
m
,
n
)
=
ck
::
type_conver
t
<
ctype
>
(
Cmn
);
}
}
}
...
...
example/19_binary_elementwise/broadcast_add_3d_am_bmnk.cpp
View file @
95ad9091
...
...
@@ -17,7 +17,8 @@ using ABDataType = F16;
using
CDataType
=
F16
;
using
EltwiseComputeDataType
=
F32
;
using
Add
=
ck
::
tensor_operation
::
binary_element_wise
::
Add
;
using
Add
=
ck
::
tensor_operation
::
binary_element_wise
::
Add
<
EltwiseComputeDataType
,
EltwiseComputeDataType
,
EltwiseComputeDataType
>
;
using
DeviceElementwiseAddInstance
=
ck
::
tensor_operation
::
device
::
DeviceBinaryElementwise
<
ABDataType
,
...
...
@@ -48,11 +49,11 @@ void host_broadcast3D_am_bmnk(HostTensorC& C,
for
(
std
::
size_t
n
=
0
;
n
<
shape
[
1
];
++
n
)
for
(
std
::
size_t
k
=
0
;
k
<
shape
[
2
];
++
k
)
{
ComputeDataType
a_val
=
static_cas
t
<
ComputeDataType
>
(
A
(
m
));
ComputeDataType
b_val
=
static_cas
t
<
ComputeDataType
>
(
B
(
m
,
n
,
k
));
ComputeDataType
a_val
=
ck
::
type_conver
t
<
ComputeDataType
>
(
A
(
m
));
ComputeDataType
b_val
=
ck
::
type_conver
t
<
ComputeDataType
>
(
B
(
m
,
n
,
k
));
ComputeDataType
c_val
=
0
;
functor
(
c_val
,
a_val
,
b_val
);
C
(
m
,
n
,
k
)
=
static_cas
t
<
ctype
>
(
c_val
);
C
(
m
,
n
,
k
)
=
ck
::
type_conver
t
<
ctype
>
(
c_val
);
}
}
...
...
example/19_binary_elementwise/elementwise_add_1d.cpp
View file @
95ad9091
/*******************************************************************************
*
* MIT License
*
* Copyright (c) 2022 Advanced Micro Devices, Inc.
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
* in the Software without restriction, including without limitation the rights
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
* copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in all
* copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*
*******************************************************************************/
#include <iostream>
#include <cstdlib>
#include "check_err.hpp"
...
...
@@ -17,7 +42,8 @@ using ABDataType = F16;
using
CDataType
=
F16
;
using
EltwiseComputeDataType
=
F32
;
using
Add
=
ck
::
tensor_operation
::
binary_element_wise
::
Add
;
using
Add
=
ck
::
tensor_operation
::
binary_element_wise
::
Add
<
EltwiseComputeDataType
,
EltwiseComputeDataType
,
EltwiseComputeDataType
>
;
using
DeviceElementwiseAddInstance
=
ck
::
tensor_operation
::
device
::
DeviceBinaryElementwise
<
ABDataType
,
...
...
@@ -43,11 +69,11 @@ void host_elementwise1D(
for
(
int
m
=
0
;
m
<
M
;
++
m
)
{
ComputeDataType
Am
=
static_cas
t
<
ComputeDataType
>
(
A
(
m
));
ComputeDataType
Bm
=
static_cas
t
<
ComputeDataType
>
(
B
(
m
));
ComputeDataType
Am
=
ck
::
type_conver
t
<
ComputeDataType
>
(
A
(
m
));
ComputeDataType
Bm
=
ck
::
type_conver
t
<
ComputeDataType
>
(
B
(
m
));
ComputeDataType
Cm
=
0
;
functor
(
Cm
,
Am
,
Bm
);
C
(
m
)
=
static_cas
t
<
ctype
>
(
Cm
);
C
(
m
)
=
ck
::
type_conver
t
<
ctype
>
(
Cm
);
}
}
...
...
example/19_binary_elementwise/elementwise_add_4d.cpp
View file @
95ad9091
/*******************************************************************************
*
* MIT License
*
* Copyright (c) 2020 Advanced Micro Devices, Inc.
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
* in the Software without restriction, including without limitation the rights
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
* copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in all
* copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*
*******************************************************************************/
#include <iostream>
#include <cstdlib>
#include "check_err.hpp"
...
...
@@ -17,7 +42,8 @@ using ABDataType = F16;
using
CDataType
=
F16
;
using
EltwiseComputeDataType
=
F32
;
using
Add
=
ck
::
tensor_operation
::
binary_element_wise
::
Add
;
using
Add
=
ck
::
tensor_operation
::
binary_element_wise
::
Add
<
EltwiseComputeDataType
,
EltwiseComputeDataType
,
EltwiseComputeDataType
>
;
using
DeviceElementwiseAddInstance
=
ck
::
tensor_operation
::
device
::
DeviceBinaryElementwise
<
ABDataType
,
...
...
@@ -49,11 +75,11 @@ void host_elementwise4D(HostTensorC& C,
for
(
std
::
size_t
h
=
0
;
h
<
shape
[
2
];
++
h
)
for
(
std
::
size_t
w
=
0
;
w
<
shape
[
3
];
++
w
)
{
ComputeDataType
a_val
=
static_cas
t
<
ComputeDataType
>
(
A
(
n
,
c
,
h
,
w
));
ComputeDataType
b_val
=
static_cas
t
<
ComputeDataType
>
(
B
(
n
,
c
,
h
,
w
));
ComputeDataType
a_val
=
ck
::
type_conver
t
<
ComputeDataType
>
(
A
(
n
,
c
,
h
,
w
));
ComputeDataType
b_val
=
ck
::
type_conver
t
<
ComputeDataType
>
(
B
(
n
,
c
,
h
,
w
));
ComputeDataType
c_val
=
0
;
functor
(
c_val
,
a_val
,
b_val
);
C
(
n
,
c
,
h
,
w
)
=
static_cas
t
<
ctype
>
(
c_val
);
C
(
n
,
c
,
h
,
w
)
=
ck
::
type_conver
t
<
ctype
>
(
c_val
);
}
}
...
...
example/21_gemm_layernorm/gemm_layernorm_xdl_fp16.cpp
View file @
95ad9091
...
...
@@ -157,8 +157,8 @@ void host_gemm_layernorm(Tensor<LayerNormOutDataType>& out_m_n,
auto
reduceSumOpInst
=
ReduceSumOp
{};
for
(
int
m
=
0
;
m
<
M
;
++
m
)
{
float
mean_acc
=
reduceSumOpInst
.
Get
ReductionZero
Val
();
float
square_mean_acc
=
reduceSumOpInst
.
Get
ReductionZero
Val
();
float
mean_acc
=
reduceSumOpInst
.
Get
Identity
Val
ue
();
float
square_mean_acc
=
reduceSumOpInst
.
Get
Identity
Val
ue
();
for
(
int
n
=
0
;
n
<
N
;
++
n
)
{
...
...
example/22_cgemm/CMakeLists.txt
0 → 100644
View file @
95ad9091
add_example_executable
(
example_cgemm_xdl_fp16 cgemm_xdl_fp16.cpp
)
example/22_cgemm/cgemm_xdl_fp16.cpp
0 → 100644
View file @
95ad9091
/*******************************************************************************
*
* MIT License
*
* Copyright (c) 2022 Advanced Micro Devices, Inc.
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
* in the Software without restriction, including without limitation the rights
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
* copies of the Software, and to permit persons to whom the Software is
* furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice shall be included in all
* copies or substantial portions of the Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
* SOFTWARE.
*
*******************************************************************************/
#include <iostream>
#include <numeric>
#include <initializer_list>
#include <cstdlib>
#include <stdlib.h>
#include <half.hpp>
#include "check_err.hpp"
#include "config.hpp"
#include "device.hpp"
#include "host_tensor.hpp"
#include "host_tensor_generator.hpp"
#include "device_tensor.hpp"
#include "device_cgemm_4gemm_xdl_cshuffle.hpp"
#include "element_wise_operation.hpp"
#include "reference_cgemm.hpp"
#include "gemm_specialization.hpp"
template
<
ck
::
index_t
...
Is
>
using
S
=
ck
::
Sequence
<
Is
...
>
;
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
using
Row
=
ck
::
tensor_layout
::
gemm
::
RowMajor
;
using
Col
=
ck
::
tensor_layout
::
gemm
::
ColumnMajor
;
using
PassThrough
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
using
ADataType
=
F16
;
using
BDataType
=
F16
;
using
CDataType
=
F16
;
using
AccDataType
=
F32
;
using
ALayout
=
ck
::
tensor_layout
::
gemm
::
RowMajor
;
using
BLayout
=
ck
::
tensor_layout
::
gemm
::
ColumnMajor
;
using
CLayout
=
ck
::
tensor_layout
::
gemm
::
RowMajor
;
static
constexpr
auto
GemmDefault
=
ck
::
tensor_operation
::
device
::
GemmSpecialization
::
Default
;
// clang-format off
using
DeviceCGemmInstance
=
ck
::
tensor_operation
::
device
::
DeviceCGemm_4Gemm_Xdl_CShuffle
<
ALayout
,
// typename ALayout
BLayout
,
// typename BLayout
CLayout
,
// typename CLayout
ADataType
,
// typename ADataType
BDataType
,
// typename BDataType
CDataType
,
// typename CDataType
AccDataType
,
// typename GemmAccDataType
CDataType
,
// typename CShuffleDataType
PassThrough
,
// typename AElementwiseOperation
PassThrough
,
// typename BElementwiseOperation
PassThrough
,
// typename CElementwiseOperation
GemmDefault
,
// GemmSpecialization GemmSpec
1
,
// index_t NumGemmKPrefetchStage
256
,
// index_t BlockSize
256
,
// index_t MPerBlock
128
,
// index_t NPerBlock
32
,
// index_t KPerBlock
8
,
// index_t AK1
8
,
// index_t BK1
32
,
// index_t MPerXDL
32
,
// index_t NPerXDL
4
,
// index_t MXdlPerWave
2
,
// index_t NXdlPerWave
S
<
4
,
64
,
1
>
,
// typename ABlockTransferThreadClusterLengths_AK0_M_AK1
S
<
1
,
0
,
2
>
,
// typename ABlockTransferThreadClusterArrangeOrder
S
<
1
,
0
,
2
>
,
// typename ABlockTransferSrcAccessOrder
2
,
// index_t ABlockTransferSrcVectorDim
8
,
// index_t ABlockTransferSrcScalarPerVector
8
,
// index_t ABlockTransferDstScalarPerVector_AK1
1
,
// index_t ABlockLdsExtraM
S
<
4
,
64
,
1
>
,
// typename BBlockTransferThreadClusterLengths_BK0_N_BK1
S
<
1
,
0
,
2
>
,
// typename BBlockTransferThreadClusterArrangeOrder
S
<
1
,
0
,
2
>
,
// typename BBlockTransferSrcAccessOrder
2
,
// index_t BBlockTransferSrcVectorDim
8
,
// index_t BBlockTransferSrcScalarPerVector
8
,
// index_t BBlockTransferDstScalarPerVector_BK1
1
,
// index_t BBlockLdsExtraN
1
,
// index_t CShuffleMXdlPerWavePerShuffle
1
,
// index_t CShuffleNXdlPerWavePerShuffle
S
<
1
,
32
,
1
,
8
>
,
// typename CShuffleBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock
8
>
;
// index_t CShuffleBlockTransferScalarPerVector_NPerBlock
// clang-format on
using
ReferenceCGemmInstance
=
ck
::
tensor_operation
::
host
::
ReferenceCGemm
<
ADataType
,
BDataType
,
CDataType
,
PassThrough
,
PassThrough
,
PassThrough
>
;
int
main
(
int
argc
,
char
*
argv
[])
{
bool
do_verification
=
true
;
int
init_method
=
1
;
bool
time_kernel
=
false
;
// CGEMM shape
ck
::
index_t
M
=
3840
;
ck
::
index_t
N
=
4096
;
ck
::
index_t
K
=
4096
;
ck
::
index_t
StrideA
=
4096
;
ck
::
index_t
StrideB
=
4096
;
ck
::
index_t
StrideC
=
4096
;
if
(
argc
==
4
)
{
do_verification
=
std
::
stoi
(
argv
[
1
]);
init_method
=
std
::
stoi
(
argv
[
2
]);
time_kernel
=
std
::
stoi
(
argv
[
3
]);
}
else
if
(
argc
==
10
)
{
do_verification
=
std
::
stoi
(
argv
[
1
]);
init_method
=
std
::
stoi
(
argv
[
2
]);
time_kernel
=
std
::
stoi
(
argv
[
3
]);
M
=
std
::
stoi
(
argv
[
4
]);
N
=
std
::
stoi
(
argv
[
5
]);
K
=
std
::
stoi
(
argv
[
6
]);
StrideA
=
std
::
stoi
(
argv
[
7
]);
StrideB
=
std
::
stoi
(
argv
[
8
]);
StrideC
=
std
::
stoi
(
argv
[
9
]);
}
else
{
printf
(
"arg1: verification (0=no, 1=yes)
\n
"
);
printf
(
"arg2: initialization (0=no init, 1=integer value, 2=decimal value)
\n
"
);
printf
(
"arg3: run kernel # of times (>1)
\n
"
);
printf
(
"arg4 to 9: M (256x), N(128x), K(32x), StrideA, StrideB, StrideC
\n
"
);
exit
(
0
);
}
auto
f_host_tensor_descriptor
=
[](
std
::
size_t
row
,
std
::
size_t
col
,
std
::
size_t
stride
,
auto
layout
)
{
if
(
std
::
is_same
<
decltype
(
layout
),
ck
::
tensor_layout
::
gemm
::
RowMajor
>::
value
)
{
return
HostTensorDescriptor
(
std
::
vector
<
std
::
size_t
>
({
row
,
col
}),
std
::
vector
<
std
::
size_t
>
({
stride
,
1
}));
}
else
{
return
HostTensorDescriptor
(
std
::
vector
<
std
::
size_t
>
({
row
,
col
}),
std
::
vector
<
std
::
size_t
>
({
1
,
stride
}));
}
};
Tensor
<
ADataType
>
a_m_k_real
(
f_host_tensor_descriptor
(
M
,
K
,
StrideA
,
ALayout
{}));
Tensor
<
ADataType
>
a_m_k_imag
(
f_host_tensor_descriptor
(
M
,
K
,
StrideA
,
ALayout
{}));
Tensor
<
BDataType
>
b_k_n_real
(
f_host_tensor_descriptor
(
K
,
N
,
StrideB
,
BLayout
{}));
Tensor
<
BDataType
>
b_k_n_imag
(
f_host_tensor_descriptor
(
K
,
N
,
StrideB
,
BLayout
{}));
Tensor
<
CDataType
>
c_m_n_real_device_result
(
f_host_tensor_descriptor
(
M
,
N
,
StrideC
,
CLayout
{}));
Tensor
<
CDataType
>
c_m_n_imag_device_result
(
f_host_tensor_descriptor
(
M
,
N
,
StrideC
,
CLayout
{}));
std
::
cout
<<
"a_m_k_real: "
<<
a_m_k_real
.
mDesc
<<
std
::
endl
;
std
::
cout
<<
"a_m_k_imag: "
<<
a_m_k_imag
.
mDesc
<<
std
::
endl
;
std
::
cout
<<
"b_k_n_real: "
<<
b_k_n_real
.
mDesc
<<
std
::
endl
;
std
::
cout
<<
"b_k_n_imag: "
<<
b_k_n_imag
.
mDesc
<<
std
::
endl
;
std
::
cout
<<
"c_m_n_real: "
<<
c_m_n_real_device_result
.
mDesc
<<
std
::
endl
;
std
::
cout
<<
"c_m_n_imag: "
<<
c_m_n_imag_device_result
.
mDesc
<<
std
::
endl
;
switch
(
init_method
)
{
case
0
:
break
;
case
1
:
a_m_k_real
.
GenerateTensorValue
(
GeneratorTensor_2
<
ADataType
>
{
-
2
,
2
});
a_m_k_imag
.
GenerateTensorValue
(
GeneratorTensor_2
<
ADataType
>
{
-
2
,
2
});
b_k_n_real
.
GenerateTensorValue
(
GeneratorTensor_2
<
BDataType
>
{
-
2
,
2
});
b_k_n_imag
.
GenerateTensorValue
(
GeneratorTensor_2
<
BDataType
>
{
-
2
,
2
});
break
;
default:
a_m_k_real
.
GenerateTensorValue
(
GeneratorTensor_3
<
ADataType
>
{
-
0.5
,
0.5
});
a_m_k_imag
.
GenerateTensorValue
(
GeneratorTensor_3
<
ADataType
>
{
-
0.5
,
0.5
});
b_k_n_real
.
GenerateTensorValue
(
GeneratorTensor_3
<
BDataType
>
{
-
0.5
,
0.5
});
b_k_n_imag
.
GenerateTensorValue
(
GeneratorTensor_3
<
BDataType
>
{
-
0.5
,
0.5
});
}
auto
cgemm
=
DeviceCGemmInstance
{};
DeviceMem
a_m_k_real_device_buf
(
sizeof
(
ADataType
)
*
a_m_k_real
.
mDesc
.
GetElementSpace
());
DeviceMem
a_m_k_imag_device_buf
(
sizeof
(
ADataType
)
*
a_m_k_imag
.
mDesc
.
GetElementSpace
());
DeviceMem
b_k_n_real_device_buf
(
sizeof
(
BDataType
)
*
b_k_n_real
.
mDesc
.
GetElementSpace
());
DeviceMem
b_k_n_imag_device_buf
(
sizeof
(
BDataType
)
*
b_k_n_imag
.
mDesc
.
GetElementSpace
());
DeviceMem
c_m_n_real_device_buf
(
sizeof
(
CDataType
)
*
c_m_n_real_device_result
.
mDesc
.
GetElementSpace
());
DeviceMem
c_m_n_imag_device_buf
(
sizeof
(
CDataType
)
*
c_m_n_imag_device_result
.
mDesc
.
GetElementSpace
());
DeviceMem
workspace_device_buf
(
cgemm
.
GetWorkspaceSize
(
M
,
N
,
K
,
StrideA
,
StrideB
,
StrideC
));
a_m_k_real_device_buf
.
ToDevice
(
a_m_k_real
.
mData
.
data
());
a_m_k_imag_device_buf
.
ToDevice
(
a_m_k_imag
.
mData
.
data
());
b_k_n_real_device_buf
.
ToDevice
(
b_k_n_real
.
mData
.
data
());
b_k_n_imag_device_buf
.
ToDevice
(
b_k_n_imag
.
mData
.
data
());
auto
a_element_op
=
PassThrough
{};
auto
b_element_op
=
PassThrough
{};
auto
c_element_op
=
PassThrough
{};
// do GEMM
auto
invoker
=
cgemm
.
MakeInvoker
();
auto
argument
=
cgemm
.
MakeArgument
(
static_cast
<
ADataType
*>
(
a_m_k_real_device_buf
.
GetDeviceBuffer
()),
static_cast
<
ADataType
*>
(
a_m_k_imag_device_buf
.
GetDeviceBuffer
()),
static_cast
<
BDataType
*>
(
b_k_n_real_device_buf
.
GetDeviceBuffer
()),
static_cast
<
BDataType
*>
(
b_k_n_imag_device_buf
.
GetDeviceBuffer
()),
static_cast
<
CDataType
*>
(
c_m_n_real_device_buf
.
GetDeviceBuffer
()),
static_cast
<
CDataType
*>
(
c_m_n_imag_device_buf
.
GetDeviceBuffer
()),
static_cast
<
CDataType
*>
(
workspace_device_buf
.
GetDeviceBuffer
()),
M
,
N
,
K
,
StrideA
,
StrideB
,
StrideC
,
a_element_op
,
b_element_op
,
c_element_op
);
if
(
!
cgemm
.
IsSupportedArgument
(
argument
))
{
throw
std
::
runtime_error
(
"wrong! device_cgemm with the specified compilation parameters does "
"not support this CGEMM problem"
);
}
float
ave_time
=
invoker
.
Run
(
argument
,
StreamConfig
{
nullptr
,
time_kernel
});
std
::
size_t
flop
=
std
::
size_t
(
8
)
*
M
*
N
*
K
;
std
::
size_t
num_btype
=
std
::
size_t
(
2
)
*
(
sizeof
(
ADataType
)
*
M
*
K
+
sizeof
(
BDataType
)
*
K
*
N
+
sizeof
(
CDataType
)
*
M
*
N
);
float
tflops
=
static_cast
<
float
>
(
flop
)
/
1.E9
/
ave_time
;
float
gb_per_sec
=
num_btype
/
1.E6
/
ave_time
;
std
::
cout
<<
"Perf: "
<<
ave_time
<<
" ms, "
<<
tflops
<<
" TFlops, "
<<
gb_per_sec
<<
" GB/s, "
<<
cgemm
.
GetTypeString
()
<<
std
::
endl
;
c_m_n_real_device_buf
.
FromDevice
(
c_m_n_real_device_result
.
mData
.
data
());
c_m_n_imag_device_buf
.
FromDevice
(
c_m_n_imag_device_result
.
mData
.
data
());
if
(
do_verification
)
{
Tensor
<
CDataType
>
c_m_n_real_host_result
(
f_host_tensor_descriptor
(
M
,
N
,
StrideC
,
CLayout
{}));
Tensor
<
CDataType
>
c_m_n_imag_host_result
(
f_host_tensor_descriptor
(
M
,
N
,
StrideC
,
CLayout
{}));
auto
ref_cgemm
=
ReferenceCGemmInstance
{};
auto
ref_invoker
=
ref_cgemm
.
MakeInvoker
();
auto
ref_argument
=
ref_cgemm
.
MakeArgument
(
a_m_k_real
,
a_m_k_imag
,
b_k_n_real
,
b_k_n_imag
,
c_m_n_real_host_result
,
c_m_n_imag_host_result
,
a_element_op
,
b_element_op
,
c_element_op
);
ref_invoker
.
Run
(
ref_argument
);
ck
::
utils
::
check_err
(
c_m_n_real_device_result
.
mData
,
c_m_n_real_host_result
.
mData
,
"Verification error: incorrect results in real part!"
,
1e-2
f
,
1e-1
f
);
ck
::
utils
::
check_err
(
c_m_n_imag_device_result
.
mData
,
c_m_n_imag_host_result
.
mData
,
"Verification error: incorrect results in imaginary part!"
,
1e-2
f
,
1e-1
f
);
}
return
0
;
}
example/CMakeLists.txt
View file @
95ad9091
...
...
@@ -48,10 +48,11 @@ add_subdirectory(11_conv2d_bwd_weight)
add_subdirectory
(
12_reduce
)
add_subdirectory
(
13_pool2d_fwd
)
add_subdirectory
(
14_gemm_xdl_requant_relu_requant
)
add_subdirectory
(
17_convnd_bwd_data_xdl
)
add_subdirectory
(
15_grouped_gemm
)
add_subdirectory
(
16_gemm_reduce
)
add_subdirectory
(
17_convnd_bwd_data_xdl
)
add_subdirectory
(
18_batched_gemm_reduce
)
add_subdirectory
(
19_binary_elementwise
)
add_subdirectory
(
20_convnd_bwd_weight_xdl
)
add_subdirectory
(
21_gemm_layernorm
)
add_subdirectory
(
22_cgemm
)
include/ck/tensor_operation/gpu/device/device_base.hpp
View file @
95ad9091
...
...
@@ -42,6 +42,8 @@ struct BaseOperator
virtual
size_t
GetWorkSpaceSize
(
const
BaseArgument
*
)
const
{
return
0
;
}
virtual
void
SetWorkSpacePointer
(
BaseArgument
*
,
void
*
)
const
{}
virtual
~
BaseOperator
()
{}
};
...
...
Prev
1
2
3
Next
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