Autotune the workspace_size_limit in conv. #40338
Merged
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
Thanks for your contribution! |
… add_kernel_metrics_tools
JamesLim-sy
force-pushed
the
add_kernel_metrics_tools
branch
from
8c6b95b
to
46ca0ac
Compare
JamesLim-sy
force-pushed
the
add_kernel_metrics_tools
branch
from
642a27e
to
7adb1b1
Compare
paddle/fluid/operators/conv_op.cc
Outdated
| @@ -583,7 +583,7 @@ void Conv3DOpMaker::Make() { | |||
| "(bool, default false) cuDNN has many algorithm to calculation " | |||
| "convolution, whether enable exhaustive search " | |||
| "for cuDNN convolution or not, default is False.") | |||
| .SetDefault(false) | |||
| .SetDefault(true) | |||
There was a problem hiding this comment.
这个默认值不能改吧？设置为true，现在默认就开启穷举搜索了。
There was a problem hiding this comment.
这里的修改是方便在Benchmark机器上跑全量模型测试，省去手动修改Benchmark中的各个脚本了。目前全量测试已经结束，会提交脚本修改回默认 false的状态。
… add_kernel_metrics_tools
…m-sy/Paddle into add_kernel_metrics_tools
…m-sy/Paddle into add_kernel_metrics_tools
JamesLim-sy
force-pushed
the
add_kernel_metrics_tools
branch
from
d5978ff
to
d68e487
Compare
JamesLim-sy
changed the title
Xreki
reviewed
Mar 28, 2022
|
|
||
| // As the container of searchAlgorithm::Find() result. | ||
| template <typename Algo_t> | ||
| struct AlgoResult { |
There was a problem hiding this comment.
Algo_t -> AlgoT
AlgoResult -> SearchResult
可以实现一个构造函数方便初始化。
| using ScalingParamType = typename platform::CudnnDataType<T>::ScalingParamType; | ||
|
|
||
| // As the basic for SearchAlgorithm struct. | ||
| template <typename conv_t> |
There was a problem hiding this comment.
原来确实是PerfT，这部分笔误了，会修改回PerfT
| }; | ||
|
|
||
| // As the container of conv relevant descriptors. | ||
| template <typename Handle_t, typename Data_t> |
| platform::FilterDescriptor wdesc; | ||
| platform::ConvolutionDescriptor cdesc; | ||
| const framework::Tensor *x, *w, *o; | ||
| Data_t cudnn_dtype; |
There was a problem hiding this comment.
platform::CudnnDataType<T>::type cudnn_dtype;
There was a problem hiding this comment.
原来是DECLARE_TYPE_FOR_GPU(dnnDataType_t, cudnnDataType_t, miopenDataType_t);
CudnnDataType和miopenDataType_t一致，根据建议修改。
There was a problem hiding this comment.
这部分修改之后，代码将隐式的模板传递，
paddle::operators::ConvArgs<T> args{&transformed_input,
&transformed_filter_channel,
&transformed_output,
strides,
padding_common,
dilations,
dtype};转变为显示传递
paddle::operators::ConvArgs<cudnnHandle_t, T> args{&transformed_input,
&transformed_filter_channel,
&transformed_output,
strides,
padding_common,
dilations,
dtype};感觉这样修改代码间接度下降了
| } | ||
|
|
||
| // template <typename algo_t> | ||
| // struct SearchAlgorithm {}; |
| *algo = result.algo; | ||
| algo_result->algo = result.algo; | ||
| algo_result->time = result.time; | ||
| algo_result->time = result.memory; | ||
| VLOG(3) << " algo: " << result.algo << ", time: " << result.time | ||
| << " ms, wksp = " << result.memory | ||
| << ", status = " << result.status; |
There was a problem hiding this comment.
打印一下workspace_bytes limit的值。
| workspace_size_limit, &algo); | ||
| ChooseAlgoByWorkspace<perf_t, algo_result_t>( | ||
| perf_results.get(), kNUM_CUDNN_FWD_ALGS, workspace_size_limit, | ||
| &result); |
| platform::dynload:: | ||
| cudnnGetConvolutionBackwardFilterAlgorithmMaxCount(args.handle, | ||
| &max_algos)); | ||
| #endif |
There was a problem hiding this comment.
可以用kNUM_CUDNN_BWD_FILTER_ALGS？
There was a problem hiding this comment.
它原来是这样的，不太清楚#if CUDNN_VERSION_MIN(7, 0, 1)之前是否有什么设置，就沿用了int max_algos = 0;
| struct KernelMetricsTool { | ||
| public: | ||
| KernelMetricsTool() {} | ||
| size_t GpuMemoryQuery() const { return GpuAvailableMemToAlloc(); } |
There was a problem hiding this comment.
该文件还会添加别的功能吗？若只是这个一个函数调用，感觉没有必要另外封装。
There was a problem hiding this comment.
嗯，核心就是调用GpuAvailableMemToAlloc()这套规则，封装中之前考虑加入GpuTimer 组件，不过既然已经实现了，这里可以删除了
| workspace_size_d = | ||
| std::max(workspace_size_d, search1::GetWorkspaceSize(args1, data_algo)); | ||
| bwd_result = search1::Find<T>(args1, exhaustive_search, deterministic, ctx); | ||
| VLOG(3) << "bwd data algo: " << bwd_result.algo << ", time " |
Xreki
force-pushed
the
add_kernel_metrics_tools
branch
from
b2e3c44
to
ea9d4b4
Compare
Xreki
changed the title
Xreki
force-pushed
the
add_kernel_metrics_tools
branch
from
de6e3e0
to
af169b6
Compare
Xreki
force-pushed
the
add_kernel_metrics_tools
branch
2 times, most recently
from
5ffbdf9
to
451a3bf
Compare
Xreki
force-pushed
the
add_kernel_metrics_tools
branch
from
451a3bf
to
ec3a2ab
Compare
lanxianghit
approved these changes
Apr 8, 2022
douch
pushed a commit
to douch/Paddle
that referenced
this pull request
Apr 10, 2022
* Using the maximum workspace_size of all alogirhms to limit the workspace size in exhaustive search mode. * Use the system cudaMalloc and cudaFree to allocate workspace during searching. * Enable switch of two kind of workspace setting methods. Co-authored-by: Liu Yiqun <liuyiqun01@baidu.com>
Xreki
added a commit
to Xreki/Paddle
that referenced
this pull request
Apr 15, 2022
* Using the maximum workspace_size of all alogirhms to limit the workspace size in exhaustive search mode. * Use the system cudaMalloc and cudaFree to allocate workspace during searching. * Enable switch of two kind of workspace setting methods. Co-authored-by: Liu Yiqun <liuyiqun01@baidu.com>
lanxianghit
pushed a commit
that referenced
this pull request
Apr 19, 2022
This pull request was closed.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
PR types
Performance optimization
PR changes
OPs
Describe
CudnnConvKernel的修改主要涵盖以下内容：通过调用
platform::GpuAvailableMemToAlloc和memory::StatGetCurrentValue，或许当前空闲系统内存、显存池内存，并此作为workspace_size_limit，动态地调整可用workspace_size的上限。FLAGS_conv_workspace_size_limit < 0时开启，FLAGS_conv_workspace_size_limit >= 0时仍使用固定的workspace_size_limit，后续将提PR改进方案并默认开启。workspace_size_limit，在搜索时会创建很小大小不同的、一次性的workspace，导致bs较小时模型训练期间显存显著增加。解决方案：搜索期间，使用cudaMalloc和cudaFree来创建workspace，避免这些workspace被显存池cache住。修改
SearchAlgorithm::Find方法返回值：定义类型SearchResult，从单独的algorithm返回值，变成返回{algorithm, time_cost, workspace_size}一组数据，同时能够在ConvKernel中获取cudnnFindXxx选择出的算法类型和时间开销，以便于提取后续的algo_profiler_result。新增
conv_base_helper.h文件：整合conv_cudnn_helper.h和conv_miopen_helper.h中的可以复用的组件，缩减无效代码规模。OP性能方面：（A100-40G机器性能）
调整前：
调整后：
SearchAlgorithm::Find返回值方面：经过测试采用
SearchResult类型的返回值后，可以在ConvKernel内得到SearchAlgorithm::Find方法获得的最优算法时间开销time_cost.模型层面
a. 在下图MLPerf ResNet50模型的第一个卷积计算中，无需设置
FLAGS_conv_workspace_size_limit，即可在FLAGS_cudnn_exhaustive_search=1时获得最佳性能。b. 通过了Benchmark平台全量模型可运行测试；
c. 通过以下常用Conv2d的PaddleClas模型中，显存打满设置的可运行测试。
该PR需要联合#41313 一起测试和验证性能，故先合进去再优化。