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
24672339
Commit
24672339
authored
May 25, 2023
by
Po-Yen, Chen
Browse files
Merge branch 'feature/integrage-karg-simplification-pr' into feature/test
parents
f2c5ca5a
853e797e
Changes
140
Hide whitespace changes
Inline
Side-by-side
Showing
20 changed files
with
556 additions
and
57 deletions
+556
-57
profiler/src/profile_max_pool3d_fwd.cpp
profiler/src/profile_max_pool3d_fwd.cpp
+168
-0
test/CMakeLists.txt
test/CMakeLists.txt
+1
-0
test/batched_gemm/CMakeLists.txt
test/batched_gemm/CMakeLists.txt
+14
-12
test/batched_gemm_gemm/CMakeLists.txt
test/batched_gemm_gemm/CMakeLists.txt
+6
-4
test/batched_gemm_reduce/CMakeLists.txt
test/batched_gemm_reduce/CMakeLists.txt
+5
-3
test/batched_gemm_softmax_gemm/CMakeLists.txt
test/batched_gemm_softmax_gemm/CMakeLists.txt
+6
-4
test/batched_gemm_softmax_gemm_permute/CMakeLists.txt
test/batched_gemm_softmax_gemm_permute/CMakeLists.txt
+15
-13
test/contraction/CMakeLists.txt
test/contraction/CMakeLists.txt
+4
-2
test/convnd_bwd_data/CMakeLists.txt
test/convnd_bwd_data/CMakeLists.txt
+4
-2
test/convnd_fwd/CMakeLists.txt
test/convnd_fwd/CMakeLists.txt
+4
-2
test/gemm_layernorm/CMakeLists.txt
test/gemm_layernorm/CMakeLists.txt
+6
-7
test/gemm_split_k/CMakeLists.txt
test/gemm_split_k/CMakeLists.txt
+5
-3
test/grouped_convnd_bwd_weight/CMakeLists.txt
test/grouped_convnd_bwd_weight/CMakeLists.txt
+4
-2
test/grouped_gemm/CMakeLists.txt
test/grouped_gemm/CMakeLists.txt
+5
-3
test/pool_fwd/CMakeLists.txt
test/pool_fwd/CMakeLists.txt
+16
-0
test/pool_fwd/test_avg_pool2d_fwd.cpp
test/pool_fwd/test_avg_pool2d_fwd.cpp
+56
-0
test/pool_fwd/test_avg_pool3d_fwd.cpp
test/pool_fwd/test_avg_pool3d_fwd.cpp
+56
-0
test/pool_fwd/test_max_pool2d_fwd.cpp
test/pool_fwd/test_max_pool2d_fwd.cpp
+75
-0
test/pool_fwd/test_max_pool3d_fwd.cpp
test/pool_fwd/test_max_pool3d_fwd.cpp
+75
-0
test/pool_fwd/test_pool_fwd_common.hpp
test/pool_fwd/test_pool_fwd_common.hpp
+31
-0
No files found.
profiler/src/profile_max_pool3d_fwd.cpp
0 → 100644
View file @
24672339
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include <iostream>
#include <vector>
#include <unordered_map>
#include "profiler/data_type_enum.hpp"
#include "profiler/profile_pool3d_fwd_impl.hpp"
#include "profiler_operation_registry.hpp"
using
ck
::
index_t
;
struct
maxPoolFwdArgParser
{
std
::
unordered_map
<
std
::
string
,
std
::
vector
<
int
>>
long_opts
=
{
{
"length"
,
{}},
{
"wsize"
,
{}},
{
"wstride"
,
{}},
{
"pad1"
,
{}},
{
"pad2"
,
{}}};
bool
parse_opt
(
int
argc
,
char
*
argv
[],
const
std
::
string
&
key
,
int
i
)
{
if
(
std
::
string
(
"--"
)
+
key
==
argv
[
i
])
{
int
pos
=
i
;
while
(
++
i
<
argc
&&
argv
[
i
][
0
]
!=
'-'
)
{}
int
end
=
i
;
for
(
int
j
=
pos
+
1
;
j
<
end
;
j
++
)
{
long_opts
[
key
].
push_back
(
std
::
stoi
(
argv
[
j
]));
}
return
true
;
}
return
false
;
}
void
operator
()(
int
argc
,
char
*
argv
[])
{
for
(
auto
&
kv
:
long_opts
)
{
for
(
int
i
=
1
;
i
<
argc
;
i
++
)
{
if
(
parse_opt
(
argc
,
argv
,
kv
.
first
,
i
))
break
;
}
}
}
};
void
print_help_max_pool3d_fwd
()
{
std
::
cout
<<
"arg1: data type (0: fp16; 1: fp32)
\n
"
<<
"arg2: verification (0: no; 1: yes)
\n
"
<<
"arg3: initialization (0: no init; 1: integer value; 2: decimal value)
\n
"
<<
"arg4: print tensor value (0: no; 1: yes)
\n
"
<<
"arg5: time kernel (0=no, 1=yes)
\n
"
<<
"arg6: return index (0=no, 1=yes)
\n
"
<<
"--length: input tensor length for NCDHW(e.g, --length 2 32 30 30 30)
\n
"
<<
"--wsize: window size for ZYX (e.g, --wsize 2 2 2)
\n
"
<<
"--wstride: window stride for DHW (e.g, --wstride 2 2 2)
\n
"
<<
"--pad1: left side of padding in DHW (e.g, --pad1 1 1 1)
\n
"
<<
"--pad2: right side of padding in DHW (e.g, --pad2 1 1 1)
\n
"
<<
"eg: ckProfiler max_pool3d_fwd 0 1 2 0 1 0 --length 2 32 30 30 30 --wsize 2 2 2 "
"--wstride 2 2 2 --pad1 1 1 1 --pad2 1 1 1"
<<
std
::
endl
;
}
int
profile_max_pool3d_fwd
(
int
argc
,
char
*
argv
[])
{
ck
::
DataTypeEnum
data_type
=
ck
::
DataTypeEnum
::
Half
;
bool
do_verification
=
true
;
int
init_method
=
0
;
bool
do_log
=
false
;
bool
time_kernel
=
true
;
bool
return_index
=
false
;
std
::
vector
<
index_t
>
in_length
=
{
2
,
32
,
30
,
30
,
30
};
std
::
vector
<
index_t
>
wsize
=
{
2
,
2
,
2
};
std
::
vector
<
index_t
>
wstride
=
{
2
,
2
,
2
};
std
::
vector
<
index_t
>
pad1
=
{
1
,
1
,
1
};
std
::
vector
<
index_t
>
pad2
=
{
1
,
1
,
1
};
if
(
argc
!=
2
&&
argc
!=
30
)
{
print_help_max_pool3d_fwd
();
return
0
;
}
else
if
(
argc
==
30
)
{
data_type
=
static_cast
<
ck
::
DataTypeEnum
>
(
std
::
stoi
(
argv
[
2
]));
do_verification
=
std
::
stoi
(
argv
[
3
]);
init_method
=
std
::
stoi
(
argv
[
4
]);
do_log
=
std
::
stoi
(
argv
[
5
]);
time_kernel
=
std
::
stoi
(
argv
[
6
]);
return_index
=
std
::
stoi
(
argv
[
7
]);
// parse the long options
maxPoolFwdArgParser
arg_parser
;
arg_parser
(
argc
,
argv
);
in_length
=
arg_parser
.
long_opts
[
"length"
];
wsize
=
arg_parser
.
long_opts
[
"wsize"
];
wstride
=
arg_parser
.
long_opts
[
"wstride"
];
pad1
=
arg_parser
.
long_opts
[
"pad1"
];
pad2
=
arg_parser
.
long_opts
[
"pad2"
];
}
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
using
I32
=
int32_t
;
constexpr
auto
ReduceOpId
=
ck
::
ReduceTensorOp
::
MAX
;
if
(
data_type
==
ck
::
DataTypeEnum
::
Half
)
{
if
(
return_index
)
ck
::
profiler
::
profile_pool3d_fwd_impl
<
F16
,
F16
,
F16
,
I32
,
ReduceOpId
,
false
,
true
>
(
do_verification
,
init_method
,
do_log
,
time_kernel
,
in_length
,
wsize
,
wstride
,
pad1
,
pad2
);
else
ck
::
profiler
::
profile_pool3d_fwd_impl
<
F16
,
F16
,
F16
,
I32
,
ReduceOpId
,
false
,
false
>
(
do_verification
,
init_method
,
do_log
,
time_kernel
,
in_length
,
wsize
,
wstride
,
pad1
,
pad2
);
}
else
if
(
data_type
==
ck
::
DataTypeEnum
::
Float
)
{
if
(
return_index
)
ck
::
profiler
::
profile_pool3d_fwd_impl
<
F32
,
F32
,
F32
,
I32
,
ReduceOpId
,
false
,
true
>
(
do_verification
,
init_method
,
do_log
,
time_kernel
,
in_length
,
wsize
,
wstride
,
pad1
,
pad2
);
else
ck
::
profiler
::
profile_pool3d_fwd_impl
<
F32
,
F32
,
F32
,
I32
,
ReduceOpId
,
false
,
false
>
(
do_verification
,
init_method
,
do_log
,
time_kernel
,
in_length
,
wsize
,
wstride
,
pad1
,
pad2
);
}
else
{
throw
std
::
runtime_error
(
"not implemented yet"
);
}
return
0
;
}
REGISTER_PROFILER_OPERATION
(
"max_pool3d_fwd"
,
"max_pool3d fwd"
,
profile_max_pool3d_fwd
);
test/CMakeLists.txt
View file @
24672339
...
...
@@ -57,6 +57,7 @@ add_subdirectory(data_type)
add_subdirectory
(
elementwise_normalization
)
add_subdirectory
(
batchnorm
)
add_subdirectory
(
contraction
)
add_subdirectory
(
pool_fwd
)
if
(
GPU_TARGETS MATCHES
"gfx1100"
)
add_subdirectory
(
wmma_op
)
endif
()
test/batched_gemm/CMakeLists.txt
View file @
24672339
add_test_executable
(
test_batched_gemm_fp16 batched_gemm_fp16.cpp
)
target_link_libraries
(
test_batched_gemm_fp16 PRIVATE utility
)
target_link_libraries
(
test_batched_gemm_fp16 PRIVATE device_batched_gemm_instance
)
if
(
GPU_TARGETS MATCHES
"gfx908"
OR GPU_TARGETS MATCHES
"gfx90a"
OR GPU_TARGETS MATCHES
"gfx940"
)
add_test_executable
(
test_batched_gemm_fp16 batched_gemm_fp16.cpp
)
target_link_libraries
(
test_batched_gemm_fp16 PRIVATE utility
)
target_link_libraries
(
test_batched_gemm_fp16 PRIVATE device_batched_gemm_instance
)
add_test_executable
(
test_batched_gemm_fp32 batched_gemm_fp32.cpp
)
target_link_libraries
(
test_batched_gemm_fp32 PRIVATE utility
)
target_link_libraries
(
test_batched_gemm_fp32 PRIVATE device_batched_gemm_instance
)
add_test_executable
(
test_batched_gemm_fp32 batched_gemm_fp32.cpp
)
target_link_libraries
(
test_batched_gemm_fp32 PRIVATE utility
)
target_link_libraries
(
test_batched_gemm_fp32 PRIVATE device_batched_gemm_instance
)
add_test_executable
(
test_batched_gemm_bf16 batched_gemm_bf16.cpp
)
target_link_libraries
(
test_batched_gemm_bf16 PRIVATE utility
)
target_link_libraries
(
test_batched_gemm_bf16 PRIVATE device_batched_gemm_instance
)
add_test_executable
(
test_batched_gemm_bf16 batched_gemm_bf16.cpp
)
target_link_libraries
(
test_batched_gemm_bf16 PRIVATE utility
)
target_link_libraries
(
test_batched_gemm_bf16 PRIVATE device_batched_gemm_instance
)
add_test_executable
(
test_batched_gemm_int8 batched_gemm_int8.cpp
)
target_link_libraries
(
test_batched_gemm_int8 PRIVATE utility
)
target_link_libraries
(
test_batched_gemm_int8 PRIVATE device_batched_gemm_instance
)
add_test_executable
(
test_batched_gemm_int8 batched_gemm_int8.cpp
)
target_link_libraries
(
test_batched_gemm_int8 PRIVATE utility
)
target_link_libraries
(
test_batched_gemm_int8 PRIVATE device_batched_gemm_instance
)
endif
()
\ No newline at end of file
test/batched_gemm_gemm/CMakeLists.txt
View file @
24672339
add_custom_target
(
test_batched_gemm_gemm
)
if
(
GPU_TARGETS MATCHES
"gfx908"
OR GPU_TARGETS MATCHES
"gfx90a"
OR GPU_TARGETS MATCHES
"gfx940"
)
add_custom_target
(
test_batched_gemm_gemm
)
add_gtest_executable
(
test_batched_gemm_gemm_fp16 test_batched_gemm_gemm_fp16.cpp
)
target_link_libraries
(
test_batched_gemm_gemm_fp16 PRIVATE utility device_batched_gemm_gemm_instance
)
add_dependencies
(
test_batched_gemm_gemm test_batched_gemm_gemm_fp16
)
\ No newline at end of file
add_gtest_executable
(
test_batched_gemm_gemm_fp16 test_batched_gemm_gemm_fp16.cpp
)
target_link_libraries
(
test_batched_gemm_gemm_fp16 PRIVATE utility device_batched_gemm_gemm_instance
)
add_dependencies
(
test_batched_gemm_gemm test_batched_gemm_gemm_fp16
)
endif
()
\ No newline at end of file
test/batched_gemm_reduce/CMakeLists.txt
View file @
24672339
add_test_executable
(
test_batched_gemm_reduce_fp16 batched_gemm_reduce_fp16.cpp
)
target_link_libraries
(
test_batched_gemm_reduce_fp16 PRIVATE utility
)
target_link_libraries
(
test_batched_gemm_reduce_fp16 PRIVATE device_batched_gemm_reduce_instance
)
if
(
GPU_TARGETS MATCHES
"gfx908"
OR GPU_TARGETS MATCHES
"gfx90a"
OR GPU_TARGETS MATCHES
"gfx940"
)
add_test_executable
(
test_batched_gemm_reduce_fp16 batched_gemm_reduce_fp16.cpp
)
target_link_libraries
(
test_batched_gemm_reduce_fp16 PRIVATE utility
)
target_link_libraries
(
test_batched_gemm_reduce_fp16 PRIVATE device_batched_gemm_reduce_instance
)
endif
()
test/batched_gemm_softmax_gemm/CMakeLists.txt
View file @
24672339
add_custom_target
(
test_batched_gemm_softmax_gemm
)
if
(
GPU_TARGETS MATCHES
"gfx908"
OR GPU_TARGETS MATCHES
"gfx90a"
OR GPU_TARGETS MATCHES
"gfx940"
)
add_custom_target
(
test_batched_gemm_softmax_gemm
)
add_gtest_executable
(
test_batched_gemm_softmax_gemm_fp16 test_batched_gemm_softmax_gemm_fp16.cpp
)
target_link_libraries
(
test_batched_gemm_softmax_gemm_fp16 PRIVATE utility device_batched_gemm_softmax_gemm_instance
)
add_dependencies
(
test_batched_gemm_softmax_gemm test_batched_gemm_softmax_gemm_fp16
)
\ No newline at end of file
add_gtest_executable
(
test_batched_gemm_softmax_gemm_fp16 test_batched_gemm_softmax_gemm_fp16.cpp
)
target_link_libraries
(
test_batched_gemm_softmax_gemm_fp16 PRIVATE utility device_batched_gemm_softmax_gemm_instance
)
add_dependencies
(
test_batched_gemm_softmax_gemm test_batched_gemm_softmax_gemm_fp16
)
endif
()
\ No newline at end of file
test/batched_gemm_softmax_gemm_permute/CMakeLists.txt
View file @
24672339
add_custom_target
(
test_batched_gemm_softmax_gemm_permute
)
if
(
GPU_TARGETS MATCHES
"gfx908"
OR GPU_TARGETS MATCHES
"gfx90a"
OR GPU_TARGETS MATCHES
"gfx940"
)
add_custom_target
(
test_batched_gemm_softmax_gemm_permute
)
add_gtest_executable
(
test_batched_gemm_softmax_gemm_permute_fp16 test_batched_gemm_softmax_gemm_permute_fp16.cpp
)
add_gtest_executable
(
test_batched_gemm_softmax_gemm_permute_bf16 test_batched_gemm_softmax_gemm_permute_bf16.cpp
)
target_link_libraries
(
test_batched_gemm_softmax_gemm_permute_fp16 PRIVATE utility device_batched_gemm_softmax_gemm_permute_instance
)
target_link_libraries
(
test_batched_gemm_softmax_gemm_permute_bf16 PRIVATE utility device_batched_gemm_softmax_gemm_permute_instance
)
add_dependencies
(
test_batched_gemm_softmax_gemm_permute test_batched_gemm_softmax_gemm_permute_fp16
)
add_dependencies
(
test_batched_gemm_softmax_gemm_permute test_batched_gemm_softmax_gemm_permute_bf16
)
add_gtest_executable
(
test_batched_gemm_softmax_gemm_permute_fp16 test_batched_gemm_softmax_gemm_permute_fp16.cpp
)
add_gtest_executable
(
test_batched_gemm_softmax_gemm_permute_bf16 test_batched_gemm_softmax_gemm_permute_bf16.cpp
)
target_link_libraries
(
test_batched_gemm_softmax_gemm_permute_fp16 PRIVATE utility device_batched_gemm_softmax_gemm_permute_instance
)
target_link_libraries
(
test_batched_gemm_softmax_gemm_permute_bf16 PRIVATE utility device_batched_gemm_softmax_gemm_permute_instance
)
add_dependencies
(
test_batched_gemm_softmax_gemm_permute test_batched_gemm_softmax_gemm_permute_fp16
)
add_dependencies
(
test_batched_gemm_softmax_gemm_permute test_batched_gemm_softmax_gemm_permute_bf16
)
add_gtest_executable
(
test_batched_gemm_bias_softmax_gemm_permute_fp16 test_batched_gemm_bias_softmax_gemm_permute_fp16.cpp
)
add_gtest_executable
(
test_batched_gemm_bias_softmax_gemm_permute_bf16 test_batched_gemm_bias_softmax_gemm_permute_bf16.cpp
)
target_link_libraries
(
test_batched_gemm_bias_softmax_gemm_permute_fp16 PRIVATE utility device_batched_gemm_softmax_gemm_permute_instance
)
target_link_libraries
(
test_batched_gemm_bias_softmax_gemm_permute_bf16 PRIVATE utility device_batched_gemm_softmax_gemm_permute_instance
)
add_dependencies
(
test_batched_gemm_softmax_gemm_permute test_batched_gemm_bias_softmax_gemm_permute_fp16
)
add_dependencies
(
test_batched_gemm_softmax_gemm_permute test_batched_gemm_bias_softmax_gemm_permute_bf16
)
\ No newline at end of file
add_gtest_executable
(
test_batched_gemm_bias_softmax_gemm_permute_fp16 test_batched_gemm_bias_softmax_gemm_permute_fp16.cpp
)
add_gtest_executable
(
test_batched_gemm_bias_softmax_gemm_permute_bf16 test_batched_gemm_bias_softmax_gemm_permute_bf16.cpp
)
target_link_libraries
(
test_batched_gemm_bias_softmax_gemm_permute_fp16 PRIVATE utility device_batched_gemm_softmax_gemm_permute_instance
)
target_link_libraries
(
test_batched_gemm_bias_softmax_gemm_permute_bf16 PRIVATE utility device_batched_gemm_softmax_gemm_permute_instance
)
add_dependencies
(
test_batched_gemm_softmax_gemm_permute test_batched_gemm_bias_softmax_gemm_permute_fp16
)
add_dependencies
(
test_batched_gemm_softmax_gemm_permute test_batched_gemm_bias_softmax_gemm_permute_bf16
)
endif
()
\ No newline at end of file
test/contraction/CMakeLists.txt
View file @
24672339
add_gtest_executable
(
test_contraction test_contraction.cpp
)
add_gtest_executable
(
test_contraction_interface test_contraction_interface.cpp
)
target_link_libraries
(
test_contraction PRIVATE utility device_contraction_bilinear_instance device_contraction_scale_instance
)
target_link_libraries
(
test_contraction_interface PRIVATE utility device_contraction_bilinear_instance device_contraction_scale_instance
)
if
(
GPU_TARGETS MATCHES
"gfx908"
OR GPU_TARGETS MATCHES
"gfx90a"
OR GPU_TARGETS MATCHES
"gfx940"
)
add_gtest_executable
(
test_contraction_interface test_contraction_interface.cpp
)
target_link_libraries
(
test_contraction_interface PRIVATE utility device_contraction_bilinear_instance device_contraction_scale_instance
)
endif
()
test/convnd_bwd_data/CMakeLists.txt
View file @
24672339
add_gtest_executable
(
test_convnd_bwd_data convnd_bwd_data.cpp
)
target_link_libraries
(
test_convnd_bwd_data PRIVATE utility device_conv1d_bwd_data_instance device_conv2d_bwd_data_instance device_conv3d_bwd_data_instance
)
if
(
GPU_TARGETS MATCHES
"gfx908"
OR GPU_TARGETS MATCHES
"gfx90a"
OR GPU_TARGETS MATCHES
"gfx940"
)
add_gtest_executable
(
test_convnd_bwd_data convnd_bwd_data.cpp
)
target_link_libraries
(
test_convnd_bwd_data PRIVATE utility device_conv1d_bwd_data_instance device_conv2d_bwd_data_instance device_conv3d_bwd_data_instance
)
endif
()
\ No newline at end of file
test/convnd_fwd/CMakeLists.txt
View file @
24672339
add_gtest_executable
(
test_convnd_fwd convnd_fwd.cpp
)
target_link_libraries
(
test_convnd_fwd PRIVATE utility device_conv2d_fwd_instance
)
if
(
GPU_TARGETS MATCHES
"gfx908"
OR GPU_TARGETS MATCHES
"gfx90a"
OR GPU_TARGETS MATCHES
"gfx940"
)
add_gtest_executable
(
test_convnd_fwd convnd_fwd.cpp
)
target_link_libraries
(
test_convnd_fwd PRIVATE utility device_conv2d_fwd_instance
)
endif
()
test/gemm_layernorm/CMakeLists.txt
View file @
24672339
add_custom_target
(
test_gemm_layernorm
)
add_gtest_executable
(
test_gemm_add_relu_add_layernorm_fp16 test_gemm_add_relu_add_layernorm_fp16.cpp
)
target_link_libraries
(
test_gemm_add_relu_add_layernorm_fp16 PRIVATE utility device_gemm_add_relu_add_layernorm_instance
)
add_dependencies
(
test_gemm_layernorm test_gemm_add_relu_add_layernorm_fp16
)
if
(
GPU_TARGETS MATCHES
"gfx908"
OR GPU_TARGETS MATCHES
"gfx90a"
OR GPU_TARGETS MATCHES
"gfx940"
)
add_custom_target
(
test_gemm_layernorm
)
add_gtest_executable
(
test_gemm_add_relu_add_layernorm_fp16 test_gemm_add_relu_add_layernorm_fp16.cpp
)
target_link_libraries
(
test_gemm_add_relu_add_layernorm_fp16 PRIVATE utility device_gemm_add_relu_add_layernorm_instance
)
add_dependencies
(
test_gemm_layernorm test_gemm_add_relu_add_layernorm_fp16
)
endif
()
test/gemm_split_k/CMakeLists.txt
View file @
24672339
add_test_executable
(
test_gemm_split_k gemm_split_k.cpp
)
target_link_libraries
(
test_gemm_split_k PRIVATE utility
)
target_link_libraries
(
test_gemm_split_k PRIVATE device_gemm_splitk_instance
)
if
(
GPU_TARGETS MATCHES
"gfx908"
OR GPU_TARGETS MATCHES
"gfx90a"
OR GPU_TARGETS MATCHES
"gfx940"
)
add_test_executable
(
test_gemm_split_k gemm_split_k.cpp
)
target_link_libraries
(
test_gemm_split_k PRIVATE utility
)
target_link_libraries
(
test_gemm_split_k PRIVATE device_gemm_splitk_instance
)
endif
()
test/grouped_convnd_bwd_weight/CMakeLists.txt
View file @
24672339
add_gtest_executable
(
test_grouped_convnd_bwd_weight grouped_convnd_bwd_weight.cpp
)
target_link_libraries
(
test_grouped_convnd_bwd_weight PRIVATE utility device_grouped_conv1d_bwd_weight_instance device_grouped_conv2d_bwd_weight_instance device_grouped_conv3d_bwd_weight_instance
)
if
(
GPU_TARGETS MATCHES
"gfx908"
OR GPU_TARGETS MATCHES
"gfx90a"
OR GPU_TARGETS MATCHES
"gfx940"
)
add_gtest_executable
(
test_grouped_convnd_bwd_weight grouped_convnd_bwd_weight.cpp
)
target_link_libraries
(
test_grouped_convnd_bwd_weight PRIVATE utility device_grouped_conv1d_bwd_weight_instance device_grouped_conv2d_bwd_weight_instance device_grouped_conv3d_bwd_weight_instance
)
endif
()
\ No newline at end of file
test/grouped_gemm/CMakeLists.txt
View file @
24672339
add_test_executable
(
test_grouped_gemm_fp16 grouped_gemm_fp16.cpp
)
target_link_libraries
(
test_grouped_gemm_fp16 PRIVATE utility
)
target_link_libraries
(
test_grouped_gemm_fp16 PRIVATE device_grouped_gemm_instance
)
if
(
GPU_TARGETS MATCHES
"gfx908"
OR GPU_TARGETS MATCHES
"gfx90a"
OR GPU_TARGETS MATCHES
"gfx940"
)
add_test_executable
(
test_grouped_gemm_fp16 grouped_gemm_fp16.cpp
)
target_link_libraries
(
test_grouped_gemm_fp16 PRIVATE utility
)
target_link_libraries
(
test_grouped_gemm_fp16 PRIVATE device_grouped_gemm_instance
)
endif
()
test/pool_fwd/CMakeLists.txt
0 → 100644
View file @
24672339
add_custom_target
(
test_pool_fwd
)
add_gtest_executable
(
test_avg_pool2d_fwd test_avg_pool2d_fwd.cpp
)
add_gtest_executable
(
test_avg_pool3d_fwd test_avg_pool3d_fwd.cpp
)
add_gtest_executable
(
test_max_pool2d_fwd test_max_pool2d_fwd.cpp
)
add_gtest_executable
(
test_max_pool3d_fwd test_max_pool3d_fwd.cpp
)
target_link_libraries
(
test_avg_pool2d_fwd PRIVATE utility device_pool_fwd_instance
)
target_link_libraries
(
test_avg_pool3d_fwd PRIVATE utility device_pool_fwd_instance
)
target_link_libraries
(
test_max_pool2d_fwd PRIVATE utility device_pool_fwd_instance
)
target_link_libraries
(
test_max_pool3d_fwd PRIVATE utility device_pool_fwd_instance
)
add_dependencies
(
test_pool_fwd test_avg_pool2d_fwd
)
add_dependencies
(
test_pool_fwd test_avg_pool3d_fwd
)
add_dependencies
(
test_pool_fwd test_max_pool2d_fwd
)
add_dependencies
(
test_pool_fwd test_max_pool3d_fwd
)
test/pool_fwd/test_avg_pool2d_fwd.cpp
0 → 100644
View file @
24672339
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "profiler/profile_pool2d_fwd_impl.hpp"
#include "test_pool_fwd_common.hpp"
template
<
typename
Tuple
>
class
TestAvgPool2dFwd
:
public
::
testing
::
Test
{
protected:
using
InDataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
OutDataType
=
std
::
tuple_element_t
<
1
,
Tuple
>
;
using
ComputeDataType
=
std
::
tuple_element_t
<
2
,
Tuple
>
;
using
IndexDataType
=
std
::
tuple_element_t
<
3
,
Tuple
>
;
std
::
vector
<
PoolingParam
>
params
;
void
Run
()
{
for
(
auto
param
:
params
)
{
bool
success
=
ck
::
profiler
::
profile_pool2d_fwd_impl
<
InDataType
,
OutDataType
,
ComputeDataType
,
IndexDataType
,
ck
::
ReduceTensorOp
::
AVG
,
false
,
false
>
(
true
,
2
,
false
,
false
,
param
.
length_
,
param
.
window_spatial_lengths_
,
param
.
window_strides_
,
param
.
input_left_pads_
,
param
.
input_right_pads_
);
EXPECT_TRUE
(
success
);
}
}
};
using
KernelTypes
=
::
testing
::
Types
<
std
::
tuple
<
F16
,
F16
,
F32
,
I32
>
,
std
::
tuple
<
F32
,
F32
,
F32
,
I32
>>
;
TYPED_TEST_SUITE
(
TestAvgPool2dFwd
,
KernelTypes
);
TYPED_TEST
(
TestAvgPool2dFwd
,
Test_Pool
)
{
// length, window_length, window_stride, left_pad, right_pad
this
->
params
=
{{{
1
,
1
,
1
,
1
},
{
1
,
1
},
{
1
,
1
},
{
0
,
0
},
{
0
,
0
}},
{{
2
,
16
,
64
,
64
},
{
64
,
64
},
{
1
,
1
},
{
0
,
0
},
{
0
,
0
}},
{{
2
,
32
,
30
,
30
},
{
2
,
2
},
{
2
,
2
},
{
1
,
1
},
{
1
,
1
}}};
this
->
Run
();
}
test/pool_fwd/test_avg_pool3d_fwd.cpp
0 → 100644
View file @
24672339
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "profiler/profile_pool3d_fwd_impl.hpp"
#include "test_pool_fwd_common.hpp"
template
<
typename
Tuple
>
class
TestAvgPool3dFwd
:
public
::
testing
::
Test
{
protected:
using
InDataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
OutDataType
=
std
::
tuple_element_t
<
1
,
Tuple
>
;
using
ComputeDataType
=
std
::
tuple_element_t
<
2
,
Tuple
>
;
using
IndexDataType
=
std
::
tuple_element_t
<
3
,
Tuple
>
;
std
::
vector
<
PoolingParam
>
params
;
void
Run
()
{
for
(
auto
param
:
params
)
{
bool
success
=
ck
::
profiler
::
profile_pool3d_fwd_impl
<
InDataType
,
OutDataType
,
ComputeDataType
,
IndexDataType
,
ck
::
ReduceTensorOp
::
AVG
,
false
,
false
>
(
true
,
2
,
false
,
false
,
param
.
length_
,
param
.
window_spatial_lengths_
,
param
.
window_strides_
,
param
.
input_left_pads_
,
param
.
input_right_pads_
);
EXPECT_TRUE
(
success
);
}
}
};
using
KernelTypes
=
::
testing
::
Types
<
std
::
tuple
<
F16
,
F16
,
F32
,
I32
>
,
std
::
tuple
<
F32
,
F32
,
F32
,
I32
>>
;
TYPED_TEST_SUITE
(
TestAvgPool3dFwd
,
KernelTypes
);
TYPED_TEST
(
TestAvgPool3dFwd
,
Test_Pool
)
{
// length, window_length, window_stride, left_pad, right_pad
this
->
params
=
{{{
1
,
1
,
1
,
1
,
1
},
{
1
,
1
,
1
},
{
1
,
1
,
1
},
{
0
,
0
,
0
},
{
0
,
0
,
0
}},
{{
2
,
16
,
64
,
64
,
64
},
{
64
,
64
,
64
},
{
1
,
1
,
1
},
{
0
,
0
,
0
},
{
0
,
0
,
0
}},
{{
2
,
32
,
30
,
30
,
30
},
{
2
,
2
,
2
},
{
2
,
2
,
2
},
{
1
,
1
,
1
},
{
1
,
1
,
1
}}};
this
->
Run
();
}
test/pool_fwd/test_max_pool2d_fwd.cpp
0 → 100644
View file @
24672339
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "profiler/profile_pool2d_fwd_impl.hpp"
#include "test_pool_fwd_common.hpp"
template
<
typename
Tuple
>
class
TestMaxPool2dFwd
:
public
::
testing
::
Test
{
protected:
using
InDataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
OutDataType
=
std
::
tuple_element_t
<
1
,
Tuple
>
;
using
ComputeDataType
=
std
::
tuple_element_t
<
2
,
Tuple
>
;
using
IndexDataType
=
std
::
tuple_element_t
<
3
,
Tuple
>
;
std
::
vector
<
PoolingParam
>
params
;
void
Run
()
{
for
(
auto
param
:
params
)
{
// max pool
bool
success
=
ck
::
profiler
::
profile_pool2d_fwd_impl
<
InDataType
,
OutDataType
,
ComputeDataType
,
IndexDataType
,
ck
::
ReduceTensorOp
::
MAX
,
false
,
false
>
(
true
,
2
,
false
,
false
,
param
.
length_
,
param
.
window_spatial_lengths_
,
param
.
window_strides_
,
param
.
input_left_pads_
,
param
.
input_right_pads_
);
EXPECT_TRUE
(
success
);
// max pool + index
success
=
ck
::
profiler
::
profile_pool2d_fwd_impl
<
InDataType
,
OutDataType
,
ComputeDataType
,
IndexDataType
,
ck
::
ReduceTensorOp
::
MAX
,
false
,
true
>
(
true
,
2
,
false
,
false
,
param
.
length_
,
param
.
window_spatial_lengths_
,
param
.
window_strides_
,
param
.
input_left_pads_
,
param
.
input_right_pads_
);
EXPECT_TRUE
(
success
);
}
}
};
using
KernelTypes
=
::
testing
::
Types
<
std
::
tuple
<
F16
,
F16
,
F16
,
I32
>
,
std
::
tuple
<
F32
,
F32
,
F32
,
I32
>>
;
TYPED_TEST_SUITE
(
TestMaxPool2dFwd
,
KernelTypes
);
TYPED_TEST
(
TestMaxPool2dFwd
,
Test_Pool
)
{
// length, window_length, window_stride, left_pad, right_pad
this
->
params
=
{{{
1
,
1
,
1
,
1
},
{
1
,
1
},
{
1
,
1
},
{
0
,
0
},
{
0
,
0
}},
{{
2
,
16
,
64
,
64
},
{
64
,
64
},
{
1
,
1
},
{
0
,
0
},
{
0
,
0
}},
{{
2
,
32
,
30
,
30
},
{
2
,
2
},
{
2
,
2
},
{
1
,
1
},
{
1
,
1
}}};
this
->
Run
();
}
test/pool_fwd/test_max_pool3d_fwd.cpp
0 → 100644
View file @
24672339
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "profiler/profile_pool3d_fwd_impl.hpp"
#include "test_pool_fwd_common.hpp"
template
<
typename
Tuple
>
class
TestMaxPool3dFwd
:
public
::
testing
::
Test
{
protected:
using
InDataType
=
std
::
tuple_element_t
<
0
,
Tuple
>
;
using
OutDataType
=
std
::
tuple_element_t
<
1
,
Tuple
>
;
using
ComputeDataType
=
std
::
tuple_element_t
<
2
,
Tuple
>
;
using
IndexDataType
=
std
::
tuple_element_t
<
3
,
Tuple
>
;
std
::
vector
<
PoolingParam
>
params
;
void
Run
()
{
for
(
auto
param
:
params
)
{
// max pool
bool
success
=
ck
::
profiler
::
profile_pool3d_fwd_impl
<
InDataType
,
OutDataType
,
ComputeDataType
,
IndexDataType
,
ck
::
ReduceTensorOp
::
MAX
,
false
,
false
>
(
true
,
2
,
false
,
false
,
param
.
length_
,
param
.
window_spatial_lengths_
,
param
.
window_strides_
,
param
.
input_left_pads_
,
param
.
input_right_pads_
);
EXPECT_TRUE
(
success
);
// max pool + index
success
=
ck
::
profiler
::
profile_pool3d_fwd_impl
<
InDataType
,
OutDataType
,
ComputeDataType
,
IndexDataType
,
ck
::
ReduceTensorOp
::
MAX
,
false
,
true
>
(
true
,
2
,
false
,
false
,
param
.
length_
,
param
.
window_spatial_lengths_
,
param
.
window_strides_
,
param
.
input_left_pads_
,
param
.
input_right_pads_
);
EXPECT_TRUE
(
success
);
}
}
};
using
KernelTypes
=
::
testing
::
Types
<
std
::
tuple
<
F16
,
F16
,
F16
,
I32
>
,
std
::
tuple
<
F32
,
F32
,
F32
,
I32
>>
;
TYPED_TEST_SUITE
(
TestMaxPool3dFwd
,
KernelTypes
);
TYPED_TEST
(
TestMaxPool3dFwd
,
Test_Pool
)
{
// length, window_length, window_stride, left_pad, right_pad
this
->
params
=
{{{
1
,
1
,
1
,
1
,
1
},
{
1
,
1
,
1
},
{
1
,
1
,
1
},
{
0
,
0
,
0
},
{
0
,
0
,
0
}},
{{
2
,
16
,
64
,
64
,
64
},
{
64
,
64
,
64
},
{
1
,
1
,
1
},
{
0
,
0
,
0
},
{
0
,
0
,
0
}},
{{
2
,
32
,
30
,
30
,
30
},
{
2
,
2
,
2
},
{
2
,
2
,
2
},
{
1
,
1
,
1
},
{
1
,
1
,
1
}}};
this
->
Run
();
}
test/pool_fwd/test_pool_fwd_common.hpp
0 → 100644
View file @
24672339
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include "gtest/gtest.h"
#include "ck/ck.hpp"
using
F16
=
ck
::
half_t
;
using
F32
=
float
;
using
I32
=
int32_t
;
using
ck
::
index_t
;
struct
PoolingParam
{
PoolingParam
(
const
std
::
vector
<
index_t
>&
length
,
const
std
::
vector
<
index_t
>&
window_spatial_lengths
,
const
std
::
vector
<
index_t
>&
window_strides
,
const
std
::
vector
<
index_t
>&
input_left_pads
,
const
std
::
vector
<
index_t
>&
input_right_pads
)
:
length_
(
length
),
window_spatial_lengths_
(
window_spatial_lengths
),
window_strides_
(
window_strides
),
input_left_pads_
(
input_left_pads
),
input_right_pads_
(
input_right_pads
)
{
}
std
::
vector
<
index_t
>
length_
;
std
::
vector
<
index_t
>
window_spatial_lengths_
;
std
::
vector
<
index_t
>
window_strides_
;
std
::
vector
<
index_t
>
input_left_pads_
;
std
::
vector
<
index_t
>
input_right_pads_
;
};
Prev
1
…
3
4
5
6
7
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