ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-12 09:16:52 +00:00
Code Issues Packages Projects Releases Wiki Activity

Modularize ckProfiler operations (#514)

Browse Source 
* Re-structure ckProfiler source files

* Rename profiler.cpp to main.cpp

* Modularize ckProfiler operations

* Add description for profiler operations

* Use longer name to avoid name collision

* Use macro to delay expansion

* Use std::move() to avoid object copying

* Prohibit users from calling dtor

* Use macro to eliminate redundant code

* Make friend function hidden

* Add missing include directive <iostream>

* Fix wrong include directives

* Remove int8 from batchnorm-forward instances since it is not needed for forward training and could fail test

Co-authored-by: Qianfeng Zhang <Qianfeng.Zhang@amd.com>
This commit is contained in:
Po Yen Chen
2022-12-02 05:15:02 +08:00
committed by GitHub
parent ad541ad6b9
commit 8784a72e23
82 changed files with 346 additions and 273 deletions
Download Patch File Download Diff File Expand all files Collapse all files

150
profiler/src/profiler.cpp
View File

@@ -1,56 +1,14 @@
// 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*[]);
int profile_batchnorm_forward(int, char*[]);
int profile_batchnorm_backward(int, char*[]);
#include "profiler_operation_registry.hpp"
static void print_helper_message()
{
// clang-format off
printf("arg1: tensor operation (gemm: GEMM\n"
" gemm_splitk: Split-K GEMM\n"
" gemm_bilinear: GEMM+Bilinear\n"
" gemm_add_add_fastgelu: GEMM+Add+Add+FastGeLU\n"
" gemm_reduce: GEMM+Reduce\n"
" gemm_bias_add_reduce: GEMM+Bias+Add+Reduce\n"
" batched_gemm: Batched GEMM\n"
" batched_gemm_gemm: Batched+GEMM+GEMM\n"
" batched_gemm_add_relu_gemm_add: Batched+GEMM+bias+gelu+GEMM+bias\n"
" batched_gemm_reduce: Batched GEMM+Reduce\n"
" grouped_gemm: Grouped GEMM\n"
" conv_fwd: Convolution Forward\n"
" conv_fwd_bias_relu: ForwardConvolution+Bias+ReLU\n"
" conv_fwd_bias_relu_add: ForwardConvolution+Bias+ReLU+Add\n"
" conv_bwd_data: Convolution Backward Data\n"
" grouped_conv_fwd: Grouped Convolution Forward\n"
" grouped_conv_bwd_weight: Grouped Convolution Backward Weight\n"
" softmax: Softmax\n"
" reduce: Reduce\n"
" bnorm_fwd: Batchnorm forward\n");
// clang-format on
std::cout << "arg1: tensor operation " << ProfilerOperationRegistry::GetInstance() << std::endl;
}
int main(int argc, char* argv[])
@@ -58,105 +16,15 @@ int main(int argc, char* argv[])
if(argc == 1)
{
print_helper_message();
return 0;
}
else if(strcmp(argv[1], "gemm") == 0)
else if(const auto operation = ProfilerOperationRegistry::GetInstance().Get(argv[1]);
operation.has_value())
{
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)
{
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 if(strcmp(argv[1], "bnorm_fwd") == 0)
{
return profile_batchnorm_forward(argc, argv);
}
else if(strcmp(argv[1], "bnorm_bwd") == 0)
{
return profile_batchnorm_backward(argc, argv);
return (*operation)(argc, argv);
}
else
{
print_helper_message();
return 0;
std::cerr << "cannot find operation: " << argv[1] << std::endl;
return EXIT_FAILURE;
}
}