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
8116d2b3
Commit
8116d2b3
authored
Nov 17, 2022
by
Po-Yen, Chen
Browse files
Modularize ckProfiler operations
parent
df021e9d
Changes
24
Hide whitespace changes
Inline
Side-by-side
Showing
4 changed files
with
99 additions
and
0 deletions
+99
-0
profiler/src/profile_reduce.cpp
profiler/src/profile_reduce.cpp
+3
-0
profiler/src/profile_softmax.cpp
profiler/src/profile_softmax.cpp
+3
-0
profiler/src/profiler.cpp
profiler/src/profiler.cpp
+47
-0
profiler/src/profiler_operation_registry.hpp
profiler/src/profiler_operation_registry.hpp
+46
-0
No files found.
profiler/src/profile_reduce.cpp
View file @
8116d2b3
...
...
@@ -15,6 +15,7 @@
#include "profiler/profile_reduce_impl.hpp"
#include "profiler/data_type_enum.hpp"
#include "profiler_operation_registry.hpp"
using
namespace
std
;
...
...
@@ -429,3 +430,5 @@ int profile_reduce(int argc, char* argv[])
return
(
0
);
};
REGISTER_PROFILER_OPERATION
(
"reduce"
,
profile_reduce
)
profiler/src/profile_softmax.cpp
View file @
8116d2b3
...
...
@@ -6,6 +6,7 @@
#include <unordered_map>
#include "profiler/profile_softmax_impl.hpp"
#include "profiler_operation_registry.hpp"
using
ck
::
index_t
;
using
ck
::
profiler
::
SoftmaxDataType
;
...
...
@@ -164,3 +165,5 @@ int profile_softmax(int argc, char* argv[])
// profile_normalization(argc, argv);
// return 0;
// }
REGISTER_PROFILER_OPERATION
(
"softmax"
,
profile_softmax
)
profiler/src/
main
.cpp
→
profiler/src/
profiler
.cpp
View file @
8116d2b3
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include <cstring>
#include <cstdlib>
#include <iostream>
int
profile_gemm
(
int
,
char
*
[]);
int
profile_gemm_splitk
(
int
,
char
*
[]);
int
profile_gemm_bilinear
(
int
,
char
*
[]);
int
profile_gemm_add_add_fastgelu
(
int
,
char
*
[]);
int
profile_gemm_reduce
(
int
,
char
*
[]);
int
profile_gemm_bias_add_reduce
(
int
,
char
*
[]);
int
profile_batched_gemm
(
int
,
char
*
[]);
int
profile_batched_gemm_gemm
(
int
,
char
*
[]);
int
profile_batched_gemm_add_relu_gemm_add
(
int
,
char
*
[]);
int
profile_batched_gemm_reduce
(
int
,
char
*
[]);
int
profile_grouped_gemm
(
int
,
char
*
[]);
int
profile_conv_fwd
(
int
,
char
*
[]);
int
profile_conv_fwd_bias_relu
(
int
,
char
*
[]);
int
profile_conv_fwd_bias_relu_add
(
int
,
char
*
[]);
int
profile_conv_bwd_data
(
int
,
char
*
[]);
int
profile_grouped_conv_fwd
(
int
,
char
*
[]);
int
profile_grouped_conv_bwd_weight
(
int
,
char
*
[]);
int
profile_softmax
(
int
,
char
*
[]);
int
profile_layernorm
(
int
,
char
*
[]);
int
profile_groupnorm
(
int
,
char
*
[]);
int
profile_reduce
(
int
,
char
*
[]);
#include "profiler_operation_registry.hpp"
static
void
print_helper_message
()
{
...
...
@@ -55,97 +36,12 @@ int main(int argc, char* argv[])
if
(
argc
==
1
)
{
print_helper_message
();
return
0
;
}
else
if
(
strcmp
(
argv
[
1
],
"gemm"
)
==
0
)
{
return
profile_gemm
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"gemm_splitk"
)
==
0
)
{
return
profile_gemm_splitk
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"gemm_bilinear"
)
==
0
)
{
return
profile_gemm_bilinear
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"gemm_add_add_fastgelu"
)
==
0
)
{
return
profile_gemm_add_add_fastgelu
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"gemm_reduce"
)
==
0
)
{
return
profile_gemm_reduce
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"gemm_bias_add_reduce"
)
==
0
)
{
return
profile_gemm_bias_add_reduce
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"batched_gemm"
)
==
0
)
{
return
profile_batched_gemm
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"batched_gemm_gemm"
)
==
0
)
{
return
profile_batched_gemm_gemm
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"batched_gemm_add_relu_gemm_add"
)
==
0
)
{
return
profile_batched_gemm_add_relu_gemm_add
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"batched_gemm_reduce"
)
==
0
)
{
return
profile_batched_gemm_reduce
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"grouped_gemm"
)
==
0
)
{
return
profile_grouped_gemm
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"conv_fwd"
)
==
0
)
else
if
(
auto
operation
=
ProfilerOperationRegistry
::
GetInstance
().
Get
(
argv
[
1
]);
operation
.
has_value
()
)
{
return
profile_conv_fwd
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"conv_fwd_bias_relu"
)
==
0
)
{
return
profile_conv_fwd_bias_relu
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"conv_fwd_bias_relu_add"
)
==
0
)
{
return
profile_conv_fwd_bias_relu_add
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"conv_bwd_data"
)
==
0
)
{
return
profile_conv_bwd_data
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"grouped_conv_fwd"
)
==
0
)
{
return
profile_grouped_conv_fwd
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"conv_bwd_weight"
)
==
0
)
{
return
profile_grouped_conv_bwd_weight
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"reduce"
)
==
0
)
{
return
profile_reduce
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"softmax"
)
==
0
)
{
return
profile_softmax
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"layernorm"
)
==
0
)
{
return
profile_layernorm
(
argc
,
argv
);
}
else
if
(
strcmp
(
argv
[
1
],
"groupnorm"
)
==
0
)
{
return
profile_groupnorm
(
argc
,
argv
);
}
else
{
print_helper_message
();
return
0
;
return
(
*
operation
)(
argc
,
argv
);
}
else
{
std
::
cerr
<<
"cannot find operation: "
<<
argv
[
1
]
<<
std
::
endl
;
return
EXIT_FAILURE
;
}
}
profiler/src/profiler_operation_registry.hpp
0 → 100644
View file @
8116d2b3
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
#include <functional>
#include <map>
#include <optional>
#include <string_view>
class
ProfilerOperationRegistry
final
{
ProfilerOperationRegistry
()
=
default
;
public:
using
Operation
=
std
::
function
<
int
(
int
,
char
*
[])
>
;
private:
std
::
unordered_map
<
std
::
string_view
,
Operation
>
operations_
;
public:
static
ProfilerOperationRegistry
&
GetInstance
()
{
static
ProfilerOperationRegistry
registry
;
return
registry
;
}
std
::
optional
<
Operation
>
Get
(
std
::
string_view
name
)
const
{
const
auto
found
=
operations_
.
find
(
name
);
if
(
found
==
end
(
operations_
))
{
return
std
::
nullopt
;
}
return
found
->
second
;
}
bool
Add
(
std
::
string_view
name
,
Operation
operation
)
{
return
operations_
.
try_emplace
(
name
,
std
::
move
(
operation
)).
second
;
}
};
#define REGISTER_PROFILER_OPERATION(name, operation) \
namespace { \
const bool result = ::ProfilerOperationRegistry::GetInstance().Add(name, operation); \
}
Prev
1
2
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