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
dacef7b5
"vscode:/vscode.git/clone" did not exist on "039b5e7ef3647d9840ac51da691e06774dc3326e"
Unverified
Commit
dacef7b5
authored
Nov 08, 2023
by
zjing14
Committed by
GitHub
Nov 08, 2023
Browse files
Merge branch 'develop' into grouped_gemm_multi_abd_fixed_nk_example
parents
70eebf22
3af8c81a
Changes
78
Show whitespace changes
Inline
Side-by-side
Showing
18 changed files
with
521 additions
and
49 deletions
+521
-49
profiler/include/profiler/profile_layernorm_fwd_impl.hpp
profiler/include/profiler/profile_layernorm_fwd_impl.hpp
+9
-9
profiler/include/profiler/profile_transpose_impl.hpp
profiler/include/profiler/profile_transpose_impl.hpp
+182
-0
profiler/src/CMakeLists.txt
profiler/src/CMakeLists.txt
+3
-3
profiler/src/profile_groupnorm_fwd.cpp
profiler/src/profile_groupnorm_fwd.cpp
+1
-1
profiler/src/profile_layernorm_fwd.cpp
profiler/src/profile_layernorm_fwd.cpp
+37
-10
profiler/src/profile_transpose.cpp
profiler/src/profile_transpose.cpp
+85
-0
test/CMakeLists.txt
test/CMakeLists.txt
+2
-1
test/normalization/CMakeLists.txt
test/normalization/CMakeLists.txt
+0
-21
test/normalization_fwd/CMakeLists.txt
test/normalization_fwd/CMakeLists.txt
+30
-0
test/normalization_fwd/test_groupnorm_fwd_fp16.cpp
test/normalization_fwd/test_groupnorm_fwd_fp16.cpp
+1
-1
test/normalization_fwd/test_groupnorm_fwd_fp32.cpp
test/normalization_fwd/test_groupnorm_fwd_fp32.cpp
+1
-1
test/normalization_fwd/test_layernorm2d_fwd_fp16.cpp
test/normalization_fwd/test_layernorm2d_fwd_fp16.cpp
+1
-1
test/normalization_fwd/test_layernorm2d_fwd_fp32.cpp
test/normalization_fwd/test_layernorm2d_fwd_fp32.cpp
+1
-1
test/normalization_fwd/test_layernorm4d_fwd_fp16.cpp
test/normalization_fwd/test_layernorm4d_fwd_fp16.cpp
+48
-0
test/transpose/CMakeLists.txt
test/transpose/CMakeLists.txt
+9
-0
test/transpose/test_transpose.cpp
test/transpose/test_transpose.cpp
+27
-0
test/transpose/test_transpose_ut_cases.inc
test/transpose/test_transpose_ut_cases.inc
+30
-0
test/transpose/test_transpose_util.hpp
test/transpose/test_transpose_util.hpp
+54
-0
No files found.
profiler/include/profiler/profile_layernorm_impl.hpp
→
profiler/include/profiler/profile_layernorm_
fwd_
impl.hpp
View file @
dacef7b5
...
@@ -6,7 +6,7 @@
...
@@ -6,7 +6,7 @@
#include <iomanip>
#include <iomanip>
#include "ck/ck.hpp"
#include "ck/ck.hpp"
#include "ck/library/tensor_operation_instance/gpu/normalization.hpp"
#include "ck/library/tensor_operation_instance/gpu/normalization
_fwd
.hpp"
#include "ck/library/utility/check_err.hpp"
#include "ck/library/utility/check_err.hpp"
#include "ck/library/utility/device_memory.hpp"
#include "ck/library/utility/device_memory.hpp"
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor.hpp"
...
@@ -94,7 +94,7 @@ bool profile_layernorm_impl(int do_verification,
...
@@ -94,7 +94,7 @@ bool profile_layernorm_impl(int do_verification,
constexpr
int
NumReduceDim
=
Rank
-
1
;
constexpr
int
NumReduceDim
=
Rank
-
1
;
// add device normalization instances
// add device normalization instances
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceNormalization
<
XDataType
,
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceNormalization
Fwd
<
XDataType
,
GammaDataType
,
GammaDataType
,
BetaDataType
,
BetaDataType
,
YDataType
,
YDataType
,
...
...
profiler/include/profiler/profile_transpose_impl.hpp
0 → 100644
View file @
dacef7b5
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include <iomanip>
#include <iostream>
#include <typeinfo>
#include "ck/ck.hpp"
#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
#include "ck/tensor_operation/gpu/device/device_elementwise.hpp"
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include "ck/tensor_operation/gpu/device/impl/device_elementwise_3d_impl.hpp"
#include "ck/library/tensor_operation_instance/gpu/transpose_3d.hpp"
#include "ck/library/utility/check_err.hpp"
#include "ck/library/utility/device_memory.hpp"
#include "ck/library/utility/host_tensor.hpp"
#include "ck/library/utility/host_tensor_generator.hpp"
#include "ck/library/utility/literals.hpp"
namespace
ck
{
namespace
profiler
{
template
<
typename
HostTensorA
,
typename
HostTensorB
,
typename
Functor
>
void
host_elementwise4D
(
HostTensorB
&
B_nchwd
,
const
HostTensorA
&
A_ncdhw
,
Functor
functor
)
{
for
(
std
::
size_t
n
=
0
;
n
<
A_ncdhw
.
mDesc
.
GetLengths
()[
0
];
++
n
)
for
(
std
::
size_t
c
=
0
;
c
<
A_ncdhw
.
mDesc
.
GetLengths
()[
1
];
++
c
)
for
(
std
::
size_t
d
=
0
;
d
<
A_ncdhw
.
mDesc
.
GetLengths
()[
2
];
++
d
)
for
(
std
::
size_t
h
=
0
;
h
<
A_ncdhw
.
mDesc
.
GetLengths
()[
3
];
++
h
)
for
(
std
::
size_t
w
=
0
;
w
<
A_ncdhw
.
mDesc
.
GetLengths
()[
4
];
++
w
)
{
auto
a_val
=
A_ncdhw
(
n
,
c
,
d
,
h
,
w
);
functor
(
B_nchwd
(
n
,
c
,
h
,
w
,
d
),
a_val
);
}
}
template
<
typename
ADataType
,
typename
BDataType
,
index_t
NumDim
>
bool
profile_transpose_impl
(
int
do_verification
,
int
init_method
,
bool
do_log
,
bool
time_kernel
,
std
::
vector
<
index_t
>
lengths
)
{
bool
pass
=
true
;
index_t
N
=
lengths
[
0
];
index_t
C
=
lengths
[
1
];
index_t
D
=
lengths
[
2
];
index_t
H
=
lengths
[
3
];
index_t
W
=
lengths
[
4
];
std
::
vector
<
ck
::
index_t
>
ncdhw
=
{
N
,
C
,
D
,
H
,
W
};
std
::
vector
<
ck
::
index_t
>
ndhwc
=
{
N
,
D
,
H
,
W
,
C
};
Tensor
<
ADataType
>
a
(
ncdhw
);
Tensor
<
BDataType
>
b
(
ndhwc
);
Tensor
<
BDataType
>
host_b
(
ndhwc
);
// a.GenerateTensorValue(GeneratorTensor_3<ADataType>{0.0, 1.0});
std
::
array
<
ck
::
index_t
,
5
>
ab_lengths
{
N
,
C
,
H
,
W
,
D
};
std
::
array
<
ck
::
index_t
,
5
>
a_strides
=
{
C
*
D
*
H
*
W
,
H
*
W
,
W
,
1
,
D
*
H
*
W
};
// N, C, D, H, W
std
::
array
<
ck
::
index_t
,
5
>
b_strides
=
{
C
*
H
*
W
*
D
,
H
*
W
*
D
,
W
*
D
,
D
,
1
};
// N, D, H, W, C
std
::
cout
<<
"A: "
<<
a
.
mDesc
<<
std
::
endl
;
std
::
cout
<<
"B: "
<<
b
.
mDesc
<<
std
::
endl
;
switch
(
init_method
)
{
case
0
:
break
;
case
1
:
a
.
GenerateTensorValue
(
GeneratorTensor_2
<
ADataType
>
{
-
1
,
2
});
break
;
default:
a
.
GenerateTensorValue
(
GeneratorTensor_3
<
ADataType
>
{
0.0
,
1.0
});
}
using
ElementOp
=
ck
::
tensor_operation
::
element_wise
::
PassThrough
;
// const auto element_op = ElementOp{};
DeviceMem
a_device_buf
(
sizeof
(
ADataType
)
*
a
.
mDesc
.
GetElementSpaceSize
());
DeviceMem
b_device_buf
(
sizeof
(
BDataType
)
*
b
.
mDesc
.
GetElementSpaceSize
());
a_device_buf
.
ToDevice
(
a
.
mData
.
data
());
std
::
array
<
const
void
*
,
1
>
input
=
{
a_device_buf
.
GetDeviceBuffer
()};
std
::
array
<
void
*
,
1
>
output
=
{
b_device_buf
.
GetDeviceBuffer
()};
using
DeviceOp
=
ck
::
tensor_operation
::
device
::
DeviceElementwise
<
ck
::
Tuple
<
ADataType
>
,
ck
::
Tuple
<
BDataType
>
,
ElementOp
,
NumDim
>
;
// get device op instances
const
auto
op_ptrs
=
ck
::
tensor_operation
::
device
::
instance
::
DeviceOperationInstanceFactory
<
DeviceOp
>::
GetInstances
();
std
::
cout
<<
"found "
<<
op_ptrs
.
size
()
<<
" instances"
<<
std
::
endl
;
if
(
do_verification
)
{
host_elementwise4D
(
host_b
,
a
,
ElementOp
{});
}
std
::
string
best_op_name
;
float
best_ave_time
=
0
;
float
best_tflops
=
0
;
float
best_gb_per_sec
=
0
;
for
(
auto
&
op_ptr
:
op_ptrs
)
{
auto
argument_ptr
=
op_ptr
->
MakeArgumentPointer
(
ab_lengths
,
{
a_strides
},
{
b_strides
},
input
,
output
,
ElementOp
{});
auto
invoker_ptr
=
op_ptr
->
MakeInvokerPointer
();
if
(
op_ptr
->
IsSupportedArgument
(
argument_ptr
.
get
()))
{
// re-init C to zero before profiling next kernel
b_device_buf
.
SetZero
();
invoker_ptr
->
Run
(
argument_ptr
.
get
(),
StreamConfig
{
nullptr
,
false
});
if
(
do_verification
)
{
b_device_buf
.
FromDevice
(
b
.
mData
.
data
());
pass
&=
ck
::
utils
::
check_err
(
b
.
mData
,
host_b
.
mData
,
"Error: Incorrect results b"
,
1e-3
,
1e-3
);
if
(
do_log
)
{
LogRangeAsType
<
float
>
(
std
::
cout
<<
"a : "
,
a
.
mData
,
","
)
<<
std
::
endl
;
LogRangeAsType
<
float
>
(
std
::
cout
<<
"b: "
,
b
.
mData
,
","
)
<<
std
::
endl
;
}
}
std
::
string
op_name
=
op_ptr
->
GetTypeString
();
float
ave_time
=
invoker_ptr
->
Run
(
argument_ptr
.
get
(),
StreamConfig
{
nullptr
,
time_kernel
});
std
::
size_t
flop
=
std
::
size_t
(
2
)
*
ncdhw
[
0
]
*
ncdhw
[
1
]
*
ncdhw
[
2
]
*
ncdhw
[
3
]
*
ncdhw
[
4
];
std
::
size_t
num_btype
=
sizeof
(
ADataType
)
*
(
ncdhw
[
0
]
*
ncdhw
[
1
]
*
ncdhw
[
2
]
*
ncdhw
[
3
]
*
ncdhw
[
4
])
+
sizeof
(
BDataType
)
*
(
ncdhw
[
0
]
*
ncdhw
[
1
]
*
ncdhw
[
2
]
*
ncdhw
[
3
]
*
ncdhw
[
4
]);
float
tflops
=
static_cast
<
float
>
(
flop
)
/
1.E9
/
ave_time
;
float
gb_per_sec
=
num_btype
/
1.E6
/
ave_time
;
std
::
cout
<<
"Perf: "
<<
std
::
setw
(
10
)
<<
ave_time
<<
" ms, "
<<
tflops
<<
" TFlops, "
<<
gb_per_sec
<<
" GB/s, "
<<
op_name
<<
std
::
endl
;
// pass = pass & ck::utils::check_err(b_device_result, b_host_result);
pass
&=
ck
::
utils
::
check_err
(
b
.
mData
,
host_b
.
mData
,
"Error: Incorrect results b"
,
1e-3
,
1e-3
);
if
(
tflops
>
best_tflops
)
{
best_op_name
=
op_name
;
best_tflops
=
tflops
;
best_ave_time
=
ave_time
;
best_gb_per_sec
=
gb_per_sec
;
}
}
else
{
std
::
cout
<<
op_ptr
->
GetTypeString
()
<<
" does not support this problem"
<<
std
::
endl
;
}
}
std
::
cout
<<
" N = "
<<
N
<<
" C = "
<<
C
<<
" D = "
<<
D
<<
" H = "
<<
H
<<
" W = "
<<
W
<<
" : "
<<
best_ave_time
<<
" ms, "
<<
best_tflops
<<
" TFlops, "
<<
best_gb_per_sec
<<
" GB/s, "
<<
best_op_name
<<
std
::
endl
;
return
pass
;
}
}
// namespace profiler
}
// namespace ck
profiler/src/CMakeLists.txt
View file @
dacef7b5
...
@@ -16,8 +16,8 @@ set(PROFILER_SOURCES
...
@@ -16,8 +16,8 @@ set(PROFILER_SOURCES
profile_grouped_conv_fwd.cpp
profile_grouped_conv_fwd.cpp
profile_grouped_conv_bwd_weight.cpp
profile_grouped_conv_bwd_weight.cpp
profile_reduce.cpp
profile_reduce.cpp
profile_groupnorm.cpp
profile_groupnorm
_fwd
.cpp
profile_layernorm.cpp
profile_layernorm
_fwd
.cpp
profile_max_pool3d_fwd.cpp
profile_max_pool3d_fwd.cpp
profile_avg_pool3d_bwd.cpp
profile_avg_pool3d_bwd.cpp
profile_max_pool3d_bwd.cpp
profile_max_pool3d_bwd.cpp
...
@@ -77,7 +77,7 @@ target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv2d_bwd_w
...
@@ -77,7 +77,7 @@ target_link_libraries(${PROFILER_EXECUTABLE} PRIVATE device_grouped_conv2d_bwd_w
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_grouped_conv3d_bwd_weight_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_grouped_conv3d_bwd_weight_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_conv2d_fwd_bias_relu_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_conv2d_fwd_bias_relu_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_conv2d_fwd_bias_relu_add_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_conv2d_fwd_bias_relu_add_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_normalization_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_normalization_
fwd_
instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_softmax_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_softmax_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_reduce_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_reduce_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_batchnorm_instance
)
target_link_libraries
(
${
PROFILER_EXECUTABLE
}
PRIVATE device_batchnorm_instance
)
...
...
profiler/src/profile_groupnorm.cpp
→
profiler/src/profile_groupnorm
_fwd
.cpp
View file @
dacef7b5
...
@@ -6,7 +6,7 @@
...
@@ -6,7 +6,7 @@
#include <unordered_map>
#include <unordered_map>
#include "profiler/data_type_enum.hpp"
#include "profiler/data_type_enum.hpp"
#include "profiler/profile_groupnorm_impl.hpp"
#include "profiler/profile_groupnorm_
fwd_
impl.hpp"
#include "profiler_operation_registry.hpp"
#include "profiler_operation_registry.hpp"
using
ck
::
index_t
;
using
ck
::
index_t
;
...
...
profiler/src/profile_layernorm.cpp
→
profiler/src/profile_layernorm
_fwd
.cpp
View file @
dacef7b5
...
@@ -6,7 +6,7 @@
...
@@ -6,7 +6,7 @@
#include <unordered_map>
#include <unordered_map>
#include "profiler/data_type_enum.hpp"
#include "profiler/data_type_enum.hpp"
#include "profiler/profile_layernorm_impl.hpp"
#include "profiler/profile_layernorm_
fwd_
impl.hpp"
#include "profiler_operation_registry.hpp"
#include "profiler_operation_registry.hpp"
using
ck
::
index_t
;
using
ck
::
index_t
;
...
@@ -78,6 +78,9 @@ int profile_layernorm(int argc, char* argv[])
...
@@ -78,6 +78,9 @@ int profile_layernorm(int argc, char* argv[])
using
F16
=
ck
::
half_t
;
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
using
F32
=
float
;
if
(
length
.
size
()
==
2
)
{
constexpr
int
rank
=
2
;
constexpr
int
rank
=
2
;
if
(
data_type
==
ck
::
DataTypeEnum
::
Half
)
if
(
data_type
==
ck
::
DataTypeEnum
::
Half
)
...
@@ -94,6 +97,30 @@ int profile_layernorm(int argc, char* argv[])
...
@@ -94,6 +97,30 @@ int profile_layernorm(int argc, char* argv[])
{
{
throw
std
::
runtime_error
(
"not implemented yet"
);
throw
std
::
runtime_error
(
"not implemented yet"
);
}
}
}
else
if
(
length
.
size
()
==
4
)
{
constexpr
int
rank
=
4
;
if
(
data_type
==
ck
::
DataTypeEnum
::
Half
)
{
ck
::
profiler
::
profile_layernorm_impl
<
F16
,
F16
,
F16
,
F32
,
F16
,
F32
,
false
,
rank
>
(
do_verification
,
init_method
,
do_log
,
time_kernel
,
length
);
}
else
if
(
data_type
==
ck
::
DataTypeEnum
::
Float
)
{
ck
::
profiler
::
profile_layernorm_impl
<
F32
,
F32
,
F32
,
F32
,
F32
,
F32
,
false
,
rank
>
(
do_verification
,
init_method
,
do_log
,
time_kernel
,
length
);
}
else
{
throw
std
::
runtime_error
(
"not implemented yet"
);
}
}
else
{
throw
std
::
runtime_error
(
"not implemented yet"
);
}
return
0
;
return
0
;
}
}
...
...
profiler/src/profile_transpose.cpp
0 → 100644
View file @
dacef7b5
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include <iostream>
#include <numeric>
#include <initializer_list>
#include <cstdlib>
#include "profiler/profile_transpose_impl.hpp"
#include "profiler_operation_registry.hpp"
enum
struct
MatrixLayout
{
NCDHW
,
// 0
NCHWD
,
// 1
};
enum
struct
DataType
{
F32_F32_F32_F32_F32
,
// 0
F16_F16_F16_F16_F16
,
// 1
};
#define OP_NAME "transpose"
#define OP_DESC "Transpose"
int
profile_transpose
(
int
argc
,
char
*
argv
[])
{
if
(
argc
!=
15
)
{
printf
(
"arg1: tensor operation ("
OP_NAME
": "
OP_DESC
")
\n
"
);
printf
(
"arg2: data type (0: fp32; 1: fp16)
\n
"
);
// printf("arg3: matrix layout (NCDHW -> NDCHW);\n");
printf
(
"arg4: verification (0: no; 1: yes)
\n
"
);
printf
(
"arg5: initialization (0: no init; 1: integer value; 2: decimal value)
\n
"
);
printf
(
"arg6: print tensor value (0: no; 1: yes)
\n
"
);
printf
(
"arg7: time kernel (0=no, 1=yes)
\n
"
);
printf
(
"arg8 to 13: N, C, D, H, W
\n
"
);
exit
(
1
);
}
const
auto
data_type
=
static_cast
<
DataType
>
(
std
::
stoi
(
argv
[
2
]));
// const auto layout = static_cast<MatrixLayout>(std::stoi(argv[3]));
const
bool
do_verification
=
std
::
stoi
(
argv
[
3
]);
const
int
init_method
=
std
::
stoi
(
argv
[
4
]);
const
bool
do_log
=
std
::
stoi
(
argv
[
5
]);
const
bool
time_kernel
=
std
::
stoi
(
argv
[
6
]);
std
::
vector
<
index_t
>
lengths
=
std
::
stoi
(
argv
[
7
]);
/**const int N = std::stoi(argv[7]);
const int C = std::stoi(argv[8]);
const int D = std::stoi(argv[9]);
const int H = std::stoi(argv[10]);
const int W = std::stoi(argv[11]);**/
using
F32
=
float
;
using
F16
=
ck
::
half_t
;
auto
profile
=
[
&
](
auto
a_type
,
auto
b_type
)
{
using
ADataType
=
decltype
(
a_type
);
using
BDataType
=
decltype
(
b_type
);
bool
pass
=
ck
::
profiler
::
profile_transpose_impl
<
ADataType
,
BDataType
>
(
do_verification
,
init_method
,
do_log
,
time_kernel
,
lengths
);
return
pass
?
0
:
1
;
};
if
(
data_type
==
GemmDataType
::
F32_F32_F32_F32_F32
)
{
return
profile
(
F32
{},
F32
{});
}
else
if
(
data_type
==
GemmDataType
::
F16_F16_F16_F16_F16
)
{
return
profile
(
F16
{},
F16
{});
}
else
{
std
::
cout
<<
"this data_type & layout is not implemented"
<<
std
::
endl
;
return
1
;
}
}
REGISTER_PROFILER_OPERATION
(
OP_NAME
,
OP_DESC
,
profile_gemm_transpose
);
test/CMakeLists.txt
View file @
dacef7b5
...
@@ -139,7 +139,7 @@ add_subdirectory(grouped_convnd_fwd)
...
@@ -139,7 +139,7 @@ add_subdirectory(grouped_convnd_fwd)
add_subdirectory
(
grouped_convnd_bwd_weight
)
add_subdirectory
(
grouped_convnd_bwd_weight
)
add_subdirectory
(
block_to_ctile_map
)
add_subdirectory
(
block_to_ctile_map
)
add_subdirectory
(
softmax
)
add_subdirectory
(
softmax
)
add_subdirectory
(
normalization
)
add_subdirectory
(
normalization
_fwd
)
add_subdirectory
(
data_type
)
add_subdirectory
(
data_type
)
add_subdirectory
(
elementwise_normalization
)
add_subdirectory
(
elementwise_normalization
)
add_subdirectory
(
batchnorm
)
add_subdirectory
(
batchnorm
)
...
@@ -148,6 +148,7 @@ add_subdirectory(pool)
...
@@ -148,6 +148,7 @@ add_subdirectory(pool)
add_subdirectory
(
batched_gemm_multi_d
)
add_subdirectory
(
batched_gemm_multi_d
)
add_subdirectory
(
grouped_convnd_bwd_data
)
add_subdirectory
(
grouped_convnd_bwd_data
)
add_subdirectory
(
conv_tensor_rearrange
)
add_subdirectory
(
conv_tensor_rearrange
)
add_subdirectory
(
transpose
)
if
(
GPU_TARGETS MATCHES
"gfx11"
)
if
(
GPU_TARGETS MATCHES
"gfx11"
)
add_subdirectory
(
wmma_op
)
add_subdirectory
(
wmma_op
)
endif
()
endif
()
test/normalization/CMakeLists.txt
deleted
100644 → 0
View file @
70eebf22
add_custom_target
(
test_normalization
)
add_gtest_executable
(
test_layernorm2d_fp32 test_layernorm2d_fp32.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_layernorm2d_fp32 PRIVATE utility device_normalization_instance
)
add_dependencies
(
test_normalization test_layernorm2d_fp32
)
endif
()
add_gtest_executable
(
test_groupnorm_fp32 test_groupnorm_fp32.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_groupnorm_fp32 PRIVATE utility device_normalization_instance
)
add_dependencies
(
test_normalization test_groupnorm_fp32
)
endif
()
add_gtest_executable
(
test_layernorm2d_fp16 test_layernorm2d_fp16.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_layernorm2d_fp16 PRIVATE utility device_normalization_instance
)
add_dependencies
(
test_normalization test_layernorm2d_fp16
)
endif
()
add_gtest_executable
(
test_groupnorm_fp16 test_groupnorm_fp16.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_groupnorm_fp16 PRIVATE utility device_normalization_instance
)
add_dependencies
(
test_normalization test_groupnorm_fp16
)
endif
()
test/normalization_fwd/CMakeLists.txt
0 → 100644
View file @
dacef7b5
add_custom_target
(
test_normalization_fwd
)
add_gtest_executable
(
test_layernorm2d_fwd_fp32 test_layernorm2d_fwd_fp32.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_layernorm2d_fwd_fp32 PRIVATE utility device_normalization_fwd_instance
)
add_dependencies
(
test_normalization_fwd test_layernorm2d_fwd_fp32
)
endif
()
add_gtest_executable
(
test_groupnorm_fwd_fp32 test_groupnorm_fwd_fp32.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_groupnorm_fwd_fp32 PRIVATE utility device_normalization_fwd_instance
)
add_dependencies
(
test_normalization_fwd test_groupnorm_fwd_fp32
)
endif
()
add_gtest_executable
(
test_layernorm2d_fwd_fp16 test_layernorm2d_fwd_fp16.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_layernorm2d_fwd_fp16 PRIVATE utility device_normalization_fwd_instance
)
add_dependencies
(
test_normalization_fwd test_layernorm2d_fwd_fp16
)
endif
()
add_gtest_executable
(
test_layernorm4d_fwd_fp16 test_layernorm4d_fwd_fp16.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_layernorm4d_fwd_fp16 PRIVATE utility device_normalization_fwd_instance
)
add_dependencies
(
test_normalization_fwd test_layernorm4d_fwd_fp16
)
endif
()
add_gtest_executable
(
test_groupnorm_fwd_fp16 test_groupnorm_fwd_fp16.cpp
)
if
(
result EQUAL 0
)
target_link_libraries
(
test_groupnorm_fwd_fp16 PRIVATE utility device_normalization_fwd_instance
)
add_dependencies
(
test_normalization_fwd test_groupnorm_fwd_fp16
)
endif
()
test/normalization/test_groupnorm_fp16.cpp
→
test/normalization
_fwd
/test_groupnorm_
fwd_
fp16.cpp
View file @
dacef7b5
...
@@ -2,7 +2,7 @@
...
@@ -2,7 +2,7 @@
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "gtest/gtest.h"
#include "profiler/profile_groupnorm_impl.hpp"
#include "profiler/profile_groupnorm_
fwd_
impl.hpp"
using
F16
=
ck
::
half_t
;
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
using
F32
=
float
;
...
...
test/normalization/test_groupnorm_fp32.cpp
→
test/normalization
_fwd
/test_groupnorm_
fwd_
fp32.cpp
View file @
dacef7b5
...
@@ -2,7 +2,7 @@
...
@@ -2,7 +2,7 @@
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "gtest/gtest.h"
#include "profiler/profile_groupnorm_impl.hpp"
#include "profiler/profile_groupnorm_
fwd_
impl.hpp"
using
F16
=
ck
::
half_t
;
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
using
F32
=
float
;
...
...
test/normalization/test_layernorm2d_fp16.cpp
→
test/normalization
_fwd
/test_layernorm2d_
fwd_
fp16.cpp
View file @
dacef7b5
...
@@ -2,7 +2,7 @@
...
@@ -2,7 +2,7 @@
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "gtest/gtest.h"
#include "profiler/profile_layernorm_impl.hpp"
#include "profiler/profile_layernorm_
fwd_
impl.hpp"
using
F16
=
ck
::
half_t
;
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
using
F32
=
float
;
...
...
test/normalization/test_layernorm2d_fp32.cpp
→
test/normalization
_fwd
/test_layernorm2d_
fwd_
fp32.cpp
View file @
dacef7b5
...
@@ -2,7 +2,7 @@
...
@@ -2,7 +2,7 @@
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "gtest/gtest.h"
#include "profiler/profile_layernorm_impl.hpp"
#include "profiler/profile_layernorm_
fwd_
impl.hpp"
using
F16
=
ck
::
half_t
;
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
using
F32
=
float
;
...
...
test/normalization_fwd/test_layernorm4d_fwd_fp16.cpp
0 → 100644
View file @
dacef7b5
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "profiler/profile_layernorm_fwd_impl.hpp"
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
using
ck
::
index_t
;
template
<
typename
Tuple
>
class
TestLayernorm4d
:
public
::
testing
::
Test
{
protected:
using
XDataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
GammaDataType
=
std
::
tuple_element_t
<
1
,
Tuple
>
;
using
BetaDataType
=
std
::
tuple_element_t
<
2
,
Tuple
>
;
using
ComputeDataType
=
std
::
tuple_element_t
<
3
,
Tuple
>
;
using
YDataType
=
std
::
tuple_element_t
<
4
,
Tuple
>
;
using
SaveMeanInvStdDataType
=
std
::
tuple_element_t
<
5
,
Tuple
>
;
void
Run
()
{
// [N, D], reduce D
std
::
vector
<
std
::
vector
<
ck
::
index_t
>>
lengths
=
{
{
1
,
1
,
1
,
1
},
{
7
,
7
,
7
,
7
},
{
256
,
16
,
16
,
8
}};
for
(
auto
length
:
lengths
)
{
bool
success
=
ck
::
profiler
::
profile_layernorm_impl
<
XDataType
,
GammaDataType
,
BetaDataType
,
ComputeDataType
,
YDataType
,
SaveMeanInvStdDataType
,
true
,
4
>
(
true
,
2
,
false
,
false
,
length
);
EXPECT_TRUE
(
success
);
}
}
};
using
KernelTypes
=
::
testing
::
Types
<
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType>
std
::
tuple
<
F16
,
F16
,
F16
,
F32
,
F16
,
F32
>>
;
TYPED_TEST_SUITE
(
TestLayernorm4d
,
KernelTypes
);
TYPED_TEST
(
TestLayernorm4d
,
Test_FP16
)
{
this
->
Run
();
}
test/transpose/CMakeLists.txt
0 → 100644
View file @
dacef7b5
list
(
APPEND gpu_list gfx908 gfx90a gfx940 gfx941 gfx942
)
set
(
target 0
)
foreach
(
gpu IN LISTS GPU_TARGETS
)
if
(
gpu IN_LIST gpu_list AND target EQUAL 0
)
add_gtest_executable
(
test_transpose test_transpose.cpp
)
target_link_libraries
(
test_transpose PRIVATE utility device_transpose_instance
)
set
(
target 1
)
endif
()
endforeach
()
test/transpose/test_transpose.cpp
0 → 100644
View file @
dacef7b5
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include <tuple>
#include "gtest/gtest.h"
#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
#include "test_transpose_util.hpp"
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
template
<
typename
Tuple
>
class
TestTranspose
:
public
::
testing
::
Test
{
};
// clang-format off
using
KernelTypes
=
::
testing
::
Types
<
std
::
tuple
<
F16
,
F16
>
,
std
::
tuple
<
F32
,
F32
>
>
;
// clang-format on
TYPED_TEST_SUITE
(
TestTranspose
,
KernelTypes
);
//#include "test_transpose_ut_cases.inc"
test/transpose/test_transpose_ut_cases.inc
0 → 100644
View file @
dacef7b5
#pragma once
TYPED_TEST
(
TestTranspose
,
Test1
)
{
// for 16, 8, 16, 32, 8
std
::
vector
<
int
>
Ms
{
1
,
2
,
3
,
4
,
5
,
6
};
std
::
vector
<
index_t
>
lengths
{
16
,
8
,
16
,
32
,
8
};
/**constexpr int N = 16;
constexpr int C = 8;
constexpr int D = 16;
constexpr int H = 32;
constexpr int W = 8;**/
this
->
Run
();
}
TYPED_TEST
(
TestTranpose
,
Test2
)
{
std
::
vector
<
int
>
Ms
{
127
,
255
,
312
,
799
,
1573
};
std
::
vector
<
index_t
>
lengths
{
16
,
8
,
16
,
32
,
16
};
/**constexpr int N = 16;
constexpr int C = 8;
constexpr int D = 16;
constexpr int H = 32;
constexpr int W = 8;**/
this
->
Run
();
}
test/transpose/test_transpose_util.hpp
0 → 100644
View file @
dacef7b5
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#pragma once
#include <string>
#include <sstream>
#include <tuple>
#include <vector>
#include <gtest/gtest.h>
#include "ck/ck.hpp"
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
#include "include/ck/utility/data_type.hpp"
#include "profiler/profile_transpose_impl.hpp"
namespace
ck
{
namespace
test
{
template
<
typename
Tuple
>
class
TestTranspose
:
public
testing
::
Test
{
using
F32
=
float
;
protected:
using
ADataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
BDataType
=
std
::
tuple_element_t
<
1
,
Tuple
>
;
public:
static
constexpr
bool
verify_
=
true
;
static
constexpr
int
init_method_
=
1
;
// decimal value initialization
static
constexpr
bool
log_
=
false
;
static
constexpr
bool
bench_
=
false
;
// measure kernel performance
std
::
vector
<
std
::
vector
<
index_t
>>
lengths_
=
{{
16
,
32
,
16
,
32
,
16
},
{
16
,
8
,
16
,
32
,
8
}};
void
Run
()
{
for
(
auto
length
:
this
->
lengths_
)
{
this
->
RunSingle
(
length
);
}
}
void
RunSingle
()
{
bool
pass
=
ck
::
profiler
::
profile_transpose_impl
<
ADataType
,
BDataType
,
5
>
(
verify_
,
init_method_
,
log_
,
bench_
,
lengths_
);
EXPECT_TRUE
(
pass
);
}
};
}
// namespace test
}
// namespace ck
Prev
1
2
3
4
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