experiments/STT-tensorflow
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

PR #36267: [ROCm] Reverting ROCm to use MIOpen Find Mode APIs (be default) for convolution

Browse Source 
Imported from GitHub PR https://github.com/tensorflow/tensorflow/pull/36267

This PR reverts ROCm to use MIOpen Find Mode APIs (be default) for convolution.  The use of MIOpen Immediate Mode API (instead of the Find Mode APIs) can be specified by the setting the env var `TF_ROCM_USE_IMMEDIATE_MODE=1`.

Almost all of the changes in this PR are within code that is specific to the ROCm platform, so this PR should not have any impact on non ROCm builds.

----------------

/cc @chsigg @whchung
Copybara import of the project:

--
5675e37e5f9b595dab45f44239cbfab222e9dcc2 by Deven Desai <deven.desai.amd@gmail.com>:

Renaming MIMIC_FIND_MODE to RETURN_BEST_ALGO_ONLY. This is being done as preparation for the implementation to re-insert calls to Find Mode API. MIMIC_FIND_MODE was a poor name for what it was doing, and would have resulted in confusion once Find Mode APIs are re-inserted. This commit also simplifies the implementation associated with RETURN_BEST_ALGO_ONLY

--
5fe0ad377dc7e333acf8aac91e3333781242fe5c by Deven Desai <deven.desai.amd@gmail.com>:

changes to fix compile time warnings in rocm_dnn.cc

--
e3dcc169353646c4b5e684b7398cf1db743079cb by Deven Desai <deven.desai.amd@gmail.com>:

Making the implementation of the Conv3D Gradient kernels consistent with the implementations of all the other Conv2D/3D kernels

--
4d4a5cede3b6e959fcc06fdb6211e4c9ef5343f5 by Deven Desai <deven.desai.amd@gmail.com>:

Updating the convolution kernel implementation(s) to ensure that the AlgorithmConfig::scratch_size field is always populated correctly before it is passed as an argument to the ThenConvolve* routine(s)

--
64ffda476af322ad804d1f5b8d7a05719e2f183c by Deven Desai <deven.desai.amd@gmail.com>:

Using the workspace memory size from the AlgorithmConfig argument, instead of calling an MIOpen API to determine it (during the call the DoPrepareForConvolution)

--
d42a76e177a26124e966c83cbbb809dbdbdcabbe by Deven Desai <deven.desai.amd@gmail.com>:

Updating the ROCm XLA Convolution Algorithm Picker implementation, to use the scratch_size that was returned in the prior call to GetMIOpenAlgorithms. Note that the code to save the scratch_size information in the new custom-call instruction (once the best conv algorithm has been determined) already exists, this commit does not change that part at all. This commit modifies how the scratch_size is determined for RunGpuConv calls that happen during the call to determine the best algorithm for a given convolution

--
416aeccbfc430c71b27cbe04a57dcd1577b34fae by Deven Desai <deven.desai.amd@gmail.com>:

Changes for the TF Convolution Kernel implementation and the Stream Executor DNN layer/api to accomodate support for Find Mode. Putting in empty placeholders in places where the Find Mode implementation will live

--
253664ce7ee59bb2ffbc2b4b3fe94963e54837c1 by Deven Desai <deven.desai.amd@gmail.com>:

Re-inserting the Find Mode Implementation. It is still disabled by default

--
30debc7b11afdbc1651c860b65cdd2fba1b9ba50 by Deven Desai <deven.desai.amd@gmail.com>:

Switching the default to Find Mode

--
b0b670e6ee2eaa6823618d4aa8858846a4cbbd89 by Deven Desai <deven.desai.amd@gmail.com>:

Disabling a subtest that fails because of bug in MIOpen Find Mode. MLOpen Issue #2379 has been filed to trach the bug.

COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/tensorflow/pull/36267 from ROCmSoftwarePlatform:google_upstream_rocm_miopen_find_mode b0b670e6ee2eaa6823618d4aa8858846a4cbbd89
PiperOrigin-RevId: 305424670
Change-Id: Ibd02cd2c43f88e619bd77e996614ded0d96d42d5
This commit is contained in:
Deven Desai 2020-04-08 00:40:43 -07:00 committed by TensorFlower Gardener
parent 74c4ae61bb
commit c6667ea3f2
16 changed files with 593 additions and 473 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

20
tensorflow/compiler/xla/service/gpu/gpu_conv_algorithm_picker.cc
View File

@ -121,11 +121,11 @@ std::vector<AlgorithmDesc> GetAlgorithms(CudnnConvKind kind,
return algorithms;
}
StatusOr<std::vector<se::dnn::ProfileResult>> GetAlgorithms(
StatusOr<std::vector<se::dnn::ProfileResult>> GetMIOpenAlgorithms(
const HloCustomCallInstruction* conv,
absl::Span<se::DeviceMemoryBase> operand_buffers,
se::DeviceMemoryBase result_buffer, se::StreamExecutor* stream_exec,
se::Stream* stream) {
ScratchAllocator* scratch_allocator, se::Stream* stream) {
std::vector<se::dnn::ProfileResult> algorithms;
TF_ASSIGN_OR_RETURN(se::dnn::ConvolutionKind kind,
@ -137,8 +137,9 @@ StatusOr<std::vector<se::dnn::ProfileResult>> GetAlgorithms(
GetGpuConvParams(conv, operand_buffers, result_buffer));
bool succ = stream_exec->GetMIOpenConvolveAlgorithms(
kind, stream, dtype, params.input_descriptor, params.filter_descriptor,
params.conv_desc, params.output_descriptor, &algorithms);
kind, dtype, stream, params.input_descriptor, params.input_buf,
params.filter_descriptor, params.filter_buf, params.output_descriptor,
params.output_buf, params.conv_desc, scratch_allocator, &algorithms);
DCHECK(succ);
return algorithms;
@ -680,9 +681,12 @@ GpuConvAlgorithmPicker::PickBestAlgorithmNoCacheRocm(
ShapeUtil::ByteSizeOf(instr->shape().tuple_shapes(0))));
initialize_buffer(result_buffer);
TF_ASSIGN_OR_RETURN(std::vector<se::dnn::ProfileResult> algorithms,
GetAlgorithms(instr, absl::MakeSpan(operand_buffers),
result_buffer, stream_exec_, stream));
ScratchAllocator scratch_allocator(device_ordinal, allocator);
TF_ASSIGN_OR_RETURN(
std::vector<se::dnn::ProfileResult> algorithms,
GetMIOpenAlgorithms(instr, absl::MakeSpan(operand_buffers), result_buffer,
stream_exec_, &scratch_allocator, stream));
std::vector<AutotuneResult> profile_results;
@ -705,7 +709,6 @@ GpuConvAlgorithmPicker::PickBestAlgorithmNoCacheRocm(
AlgorithmToString(alg)),
2);
ScratchAllocator scratch_allocator(device_ordinal, allocator);
se::dnn::ProfileResult profile_result;
VLOG(3) << "Trying algorithm " << AlgorithmToString(alg) << " for "
<< instr->ToString();
@ -714,6 +717,7 @@ GpuConvAlgorithmPicker::PickBestAlgorithmNoCacheRocm(
RunConvOptions options;
options.profile_result = &profile_result;
options.algo_override = alg;
options.scratch_size_override = miopen_alg.scratch_size();
Status launch_status =
RunGpuConv(instr, absl::MakeSpan(operand_buffers), result_buffer,
&scratch_allocator, stream, options);

3
tensorflow/compiler/xla/service/gpu/gpu_conv_runner.cc
View File

@ -225,6 +225,9 @@ Status RunGpuConvImpl(const GpuConvParams& params,
if (options.algo_override.has_value()) {
algorithm = AlgorithmConfig(*options.algo_override);
if (options.scratch_size_override.has_value()) {
algorithm.set_scratch_size(*options.scratch_size_override);
}
}
Status run_status = RunGpuConvInternalImpl<ElementType, BiasType, OutputType>(

3
tensorflow/compiler/xla/service/gpu/gpu_conv_runner.h
View File

@ -35,6 +35,9 @@ struct RunConvOptions {
// Use this algorithm, instead of the one from the instruction.
absl::optional<se::dnn::AlgorithmDesc> algo_override;
// Use this scratch_bytes size, instead of the one from the instruction.
absl::optional<size_t> scratch_size_override;
};
// Implementation struct exposed for debugging and log analysis.

29
tensorflow/core/kernels/conv_grad_filter_ops.cc
View File

@ -1040,16 +1040,21 @@ void LaunchConv2DBackpropFilterOp<Eigen::GpuDevice, T>::operator()(
}
}
#elif TENSORFLOW_USE_ROCM
DnnScratchAllocator scratch_allocator(ConvolveBackwardFilterScratchSize,
ctx);
std::vector<ProfileResult> algorithms;
OP_REQUIRES(ctx,
stream->parent()->GetMIOpenConvolveAlgorithms(
se::dnn::ConvolutionKind::BACKWARD_FILTER, stream,
se::dnn::ToDataType<T>::value, input_desc, filter_desc,
conv_desc, output_desc, &algorithms),
errors::Unknown(
"Failed to get convolution algorithm. This is probably "
"because MIOpen failed to initialize, so try looking to "
"see if a warning log message was printed above."));
OP_REQUIRES(
ctx,
stream->parent()->GetMIOpenConvolveAlgorithms(
se::dnn::ConvolutionKind::BACKWARD_FILTER,
se::dnn::ToDataType<T>::value, stream, input_desc, input_ptr,
filter_desc, filter_backprop_ptr, output_desc, out_backprop_ptr,
conv_desc, &scratch_allocator, &algorithms),
errors::Unknown(
"Failed to get convolution algorithm. This is probably "
"because MIOpen failed to initialize, so try looking to "
"see if a warning log message was printed above."));
std::vector<tensorflow::AutotuneResult> results;
if (algorithms.size() == 1) {
@ -1067,8 +1072,6 @@ void LaunchConv2DBackpropFilterOp<Eigen::GpuDevice, T>::operator()(
} else {
for (auto miopen_algorithm : algorithms) {
auto profile_algorithm = miopen_algorithm.algorithm();
DnnScratchAllocator scratch_allocator(ConvolveBackwardFilterScratchSize,
ctx);
ProfileResult profile_result;
bool miopen_launch_status = true;
miopen_launch_status =
@ -1076,7 +1079,9 @@ void LaunchConv2DBackpropFilterOp<Eigen::GpuDevice, T>::operator()(
->ThenConvolveBackwardFilterWithAlgorithm(
input_desc, input_ptr, output_desc, out_backprop_ptr,
conv_desc, filter_desc, &filter_backprop_ptr,
&scratch_allocator, AlgorithmConfig(profile_algorithm),
&scratch_allocator,
AlgorithmConfig(profile_algorithm,
miopen_algorithm.scratch_size()),
&profile_result)
.ok();

27
tensorflow/core/kernels/conv_grad_input_ops.cc
View File

@ -1202,16 +1202,19 @@ void LaunchConv2DBackpropInputOp<GPUDevice, T>::operator()(
}
}
#elif TENSORFLOW_USE_ROCM
DnnScratchAllocator scratch_allocator(ConvolveBackwardDataScratchSize, ctx);
std::vector<ProfileResult> algorithms;
OP_REQUIRES(ctx,
stream->parent()->GetMIOpenConvolveAlgorithms(
se::dnn::ConvolutionKind::BACKWARD_DATA, stream,
se::dnn::ToDataType<T>::value, input_desc, filter_desc,
conv_desc, output_desc, &algorithms),
errors::Unknown(
"Failed to get convolution algorithm. This is probably "
"because MIOpen failed to initialize, so try looking to "
"see if a warning log message was printed above."));
OP_REQUIRES(
ctx,
stream->parent()->GetMIOpenConvolveAlgorithms(
se::dnn::ConvolutionKind::BACKWARD_DATA,
se::dnn::ToDataType<T>::value, stream, input_desc, in_backprop_ptr,
filter_desc, filter_ptr, output_desc, out_backprop_ptr, conv_desc,
&scratch_allocator, &algorithms),
errors::Unknown(
"Failed to get convolution algorithm. This is probably "
"because MIOpen failed to initialize, so try looking to "
"see if a warning log message was printed above."));
std::vector<tensorflow::AutotuneResult> results;
if (algorithms.size() == 1) {
@ -1229,8 +1232,6 @@ void LaunchConv2DBackpropInputOp<GPUDevice, T>::operator()(
} else {
for (auto miopen_algorithm : algorithms) {
auto profile_algorithm = miopen_algorithm.algorithm();
DnnScratchAllocator scratch_allocator(ConvolveBackwardDataScratchSize,
ctx);
ProfileResult profile_result;
bool miopen_launch_status = true;
miopen_launch_status =
@ -1238,7 +1239,9 @@ void LaunchConv2DBackpropInputOp<GPUDevice, T>::operator()(
->ThenConvolveBackwardDataWithAlgorithm(
filter_desc, filter_ptr, output_desc, out_backprop_ptr,
conv_desc, input_desc, &in_backprop_ptr, &scratch_allocator,
AlgorithmConfig(profile_algorithm), &profile_result)
AlgorithmConfig(profile_algorithm,
miopen_algorithm.scratch_size()),
&profile_result)
.ok();
if (miopen_launch_status && profile_result.is_valid()) {

153
tensorflow/core/kernels/conv_grad_ops_3d.cc
View File

@ -1389,8 +1389,7 @@ class Conv3DBackpropInputOp<GPUDevice, T> : public OpKernel {
conv_parameters.ShouldIncludeWinogradNonfusedAlgo<T>(
stream->parent()),
&algorithms));
ProfileResult best_result;
ProfileResult best_result_no_scratch;
std::vector<tensorflow::AutotuneResult> results;
for (const auto& profile_algorithm : algorithms) {
// TODO(zhengxq): profile each algorithm multiple times to better
@ -1427,15 +1426,6 @@ class Conv3DBackpropInputOp<GPUDevice, T> : public OpKernel {
*result.mutable_run_time() = proto_utils::ToDurationProto(
absl::Milliseconds(profile_result.elapsed_time_in_ms()));
if (profile_result.elapsed_time_in_ms() <
best_result.elapsed_time_in_ms()) {
best_result = profile_result;
}
if (scratch_allocator.TotalByteSize() == 0 &&
profile_result.elapsed_time_in_ms() <
best_result_no_scratch.elapsed_time_in_ms()) {
best_result_no_scratch = profile_result;
}
// TODO(george): they don't do results at all??
CheckRedzones(rz_scratch_allocator, &result);
CheckRedzones(rz_allocator, &result);
@ -1443,25 +1433,26 @@ class Conv3DBackpropInputOp<GPUDevice, T> : public OpKernel {
}
}
#elif TENSORFLOW_USE_ROCM
DnnScratchAllocator scratch_allocator(ConvolveBackwardDataScratchSize,
context);
std::vector<ProfileResult> algorithms;
CHECK(stream->parent()->GetMIOpenConvolveAlgorithms(
se::dnn::ConvolutionKind::BACKWARD_DATA, stream,
se::dnn::ToDataType<T>::value, input_desc, filter_desc, conv_desc,
output_desc, &algorithms));
ProfileResult best_result;
ProfileResult best_result_no_scratch;
se::dnn::ConvolutionKind::BACKWARD_DATA,
se::dnn::ToDataType<T>::value, stream, input_desc, in_backprop_ptr,
filter_desc, filter_ptr, output_desc, out_backprop_ptr, conv_desc,
&scratch_allocator, &algorithms));
std::vector<tensorflow::AutotuneResult> results;
for (auto miopen_algorithm : algorithms) {
auto profile_algorithm = miopen_algorithm.algorithm();
DnnScratchAllocator scratch_allocator(ConvolveBackwardDataScratchSize,
context);
ProfileResult profile_result;
bool miopen_launch_status =
stream
->ThenConvolveBackwardDataWithAlgorithm(
filter_desc, filter_ptr, output_desc, out_backprop_ptr,
conv_desc, input_desc, &in_backprop_ptr, &scratch_allocator,
AlgorithmConfig(profile_algorithm), &profile_result)
AlgorithmConfig(profile_algorithm,
miopen_algorithm.scratch_size()),
&profile_result)
.ok();
if (miopen_launch_status) {
if (profile_result.is_valid()) {
@ -1473,16 +1464,6 @@ class Conv3DBackpropInputOp<GPUDevice, T> : public OpKernel {
result.set_scratch_bytes(scratch_allocator.TotalByteSize());
*result.mutable_run_time() = proto_utils::ToDurationProto(
absl::Milliseconds(profile_result.elapsed_time_in_ms()));
if (profile_result.elapsed_time_in_ms() <
best_result.elapsed_time_in_ms()) {
best_result = profile_result;
}
if (scratch_allocator.TotalByteSize() == 0 &&
profile_result.elapsed_time_in_ms() <
best_result_no_scratch.elapsed_time_in_ms()) {
best_result_no_scratch = profile_result;
}
}
}
}
@ -1492,16 +1473,8 @@ class Conv3DBackpropInputOp<GPUDevice, T> : public OpKernel {
filter_ptr, out_backprop_ptr, input_desc,
filter_desc, output_desc, conv_desc,
stream->parent(), results);
OP_REQUIRES(context,
best_result.is_valid() || best_result_no_scratch.is_valid(),
errors::NotFound("No algorithm worked!"));
if (best_result.is_valid()) {
algorithm_config.set_algorithm(best_result.algorithm());
}
if (best_result_no_scratch.is_valid()) {
algorithm_config.set_algorithm_no_scratch(
best_result_no_scratch.algorithm());
}
OP_REQUIRES_OK(context,
BestCudnnConvAlgorithm(results, &algorithm_config));
AutoTuneConv3dBwdData::GetInstance()->Insert(conv_parameters,
algorithm_config);
}
@ -1878,8 +1851,8 @@ class Conv3DBackpropFilterOp<GPUDevice, T> : public OpKernel {
conv_parameters.ShouldIncludeWinogradNonfusedAlgo<T>(
stream->parent()),
&algorithms));
ProfileResult best_result;
ProfileResult best_result_no_scratch;
std::vector<tensorflow::AutotuneResult> results;
for (const auto& profile_algorithm : algorithms) {
// TODO(zhengxq): profile each algorithm multiple times to better
// accuracy.
@ -1896,68 +1869,62 @@ class Conv3DBackpropFilterOp<GPUDevice, T> : public OpKernel {
.ok();
if (cudnn_launch_status) {
if (profile_result.is_valid()) {
if (profile_result.elapsed_time_in_ms() <
best_result.elapsed_time_in_ms()) {
best_result = profile_result;
}
if (scratch_allocator.TotalByteSize() == 0 &&
profile_result.elapsed_time_in_ms() <
best_result_no_scratch.elapsed_time_in_ms()) {
best_result_no_scratch = profile_result;
}
results.emplace_back();
auto& result = results.back();
result.mutable_conv()->set_algorithm(profile_algorithm.algo_id());
result.mutable_conv()->set_tensor_ops_enabled(
profile_algorithm.tensor_ops_enabled());
result.set_scratch_bytes(scratch_allocator.TotalByteSize());
*result.mutable_run_time() = proto_utils::ToDurationProto(
absl::Milliseconds(profile_result.elapsed_time_in_ms()));
}
}
}
#elif TENSORFLOW_USE_ROCM
DnnScratchAllocator scratch_allocator(ConvolveBackwardFilterScratchSize,
context);
std::vector<ProfileResult> algorithms;
CHECK(stream->parent()->GetMIOpenConvolveAlgorithms(
se::dnn::ConvolutionKind::BACKWARD_FILTER, stream,
se::dnn::ToDataType<T>::value, input_desc, filter_desc, conv_desc,
output_desc, &algorithms));
ProfileResult best_result;
ProfileResult best_result_no_scratch;
if (algorithms.size() == 1) {
best_result = algorithms[0];
} else {
for (auto miopen_algorithm : algorithms) {
auto profile_algorithm = miopen_algorithm.algorithm();
DnnScratchAllocator scratch_allocator(
ConvolveBackwardFilterScratchSize, context);
ProfileResult profile_result;
bool cudnn_launch_status =
stream
->ThenConvolveBackwardFilterWithAlgorithm(
input_desc, input_ptr, output_desc, out_backprop_ptr,
conv_desc, filter_desc, &filter_backprop_ptr,
&scratch_allocator, AlgorithmConfig(profile_algorithm),
&profile_result)
.ok();
if (cudnn_launch_status) {
if (profile_result.is_valid()) {
if (profile_result.elapsed_time_in_ms() <
best_result.elapsed_time_in_ms()) {
best_result = profile_result;
}
if (scratch_allocator.TotalByteSize() == 0 &&
profile_result.elapsed_time_in_ms() <
best_result_no_scratch.elapsed_time_in_ms()) {
best_result_no_scratch = profile_result;
}
}
se::dnn::ConvolutionKind::BACKWARD_FILTER,
se::dnn::ToDataType<T>::value, stream, input_desc, input_ptr,
filter_desc, filter_backprop_ptr, output_desc, out_backprop_ptr,
conv_desc, &scratch_allocator, &algorithms));
std::vector<tensorflow::AutotuneResult> results;
for (auto miopen_algorithm : algorithms) {
auto profile_algorithm = miopen_algorithm.algorithm();
ProfileResult profile_result;
bool cudnn_launch_status =
stream
->ThenConvolveBackwardFilterWithAlgorithm(
input_desc, input_ptr, output_desc, out_backprop_ptr,
conv_desc, filter_desc, &filter_backprop_ptr,
&scratch_allocator,
AlgorithmConfig(profile_algorithm,
miopen_algorithm.scratch_size()),
&profile_result)
.ok();
if (cudnn_launch_status) {
if (profile_result.is_valid()) {
results.emplace_back();
auto& result = results.back();
result.mutable_conv()->set_algorithm(profile_algorithm.algo_id());
result.mutable_conv()->set_tensor_ops_enabled(
profile_algorithm.tensor_ops_enabled());
result.set_scratch_bytes(scratch_allocator.TotalByteSize());
*result.mutable_run_time() = proto_utils::ToDurationProto(
absl::Milliseconds(profile_result.elapsed_time_in_ms()));
}
}
}
#endif
OP_REQUIRES(context,
best_result.is_valid() || best_result_no_scratch.is_valid(),
errors::NotFound("No algorithm worked!"));
if (best_result.is_valid()) {
algorithm_config.set_algorithm(best_result.algorithm());
}
if (best_result_no_scratch.is_valid()) {
algorithm_config.set_algorithm_no_scratch(
best_result_no_scratch.algorithm());
}
LogConvAutotuneResults(se::dnn::ConvolutionKind::BACKWARD_FILTER,
se::dnn::ToDataType<T>::value, input_ptr,
filter_backprop_ptr, out_backprop_ptr, input_desc,
filter_desc, output_desc, conv_desc,
stream->parent(), results);
OP_REQUIRES_OK(context,
BestCudnnConvAlgorithm(results, &algorithm_config));
AutoTuneConv3dBwdFilter::GetInstance()->Insert(conv_parameters,
algorithm_config);
}

26
tensorflow/core/kernels/conv_ops.cc
View File

@ -1052,16 +1052,19 @@ void LaunchConv2DOp<GPUDevice, T>::operator()(
}
#elif TENSORFLOW_USE_ROCM
DnnScratchAllocator scratch_allocator(ConvolveScratchSize, ctx);
std::vector<ProfileResult> algorithms;
OP_REQUIRES(ctx,
stream->parent()->GetMIOpenConvolveAlgorithms(
se::dnn::ConvolutionKind::FORWARD, stream,
se::dnn::ToDataType<T>::value, input_desc, filter_desc,
conv_desc, output_desc, &algorithms),
errors::Unknown(
"Failed to get convolution algorithm. This is probably "
"because MIOpen failed to initialize, so try looking to "
"see if a warning log message was printed above."));
OP_REQUIRES(
ctx,
stream->parent()->GetMIOpenConvolveAlgorithms(
se::dnn::ConvolutionKind::FORWARD, se::dnn::ToDataType<T>::value,
stream, input_desc, input_ptr, filter_desc, filter_ptr, output_desc,
output_ptr, conv_desc, &scratch_allocator, &algorithms),
errors::Unknown(
"Failed to get convolution algorithm. This is probably "
"because MIOpen failed to initialize, so try looking to "
"see if a warning log message was printed above."));
se::DeviceMemory<T> output_tensor = output_ptr;
std::vector<tensorflow::AutotuneResult> results;
@ -1080,7 +1083,6 @@ void LaunchConv2DOp<GPUDevice, T>::operator()(
} else {
for (auto miopen_algorithm : algorithms) {
auto profile_algorithm = miopen_algorithm.algorithm();
DnnScratchAllocator scratch_allocator(ConvolveScratchSize, ctx);
ProfileResult profile_result;
bool miopen_launch_status = false;
miopen_launch_status =
@ -1088,7 +1090,9 @@ void LaunchConv2DOp<GPUDevice, T>::operator()(
->ThenConvolveWithAlgorithm(
input_desc, input_ptr, filter_desc, filter_ptr, conv_desc,
output_desc, &output_ptr, &scratch_allocator,
AlgorithmConfig(profile_algorithm), &profile_result)
AlgorithmConfig(profile_algorithm,
miopen_algorithm.scratch_size()),
&profile_result)
.ok();
if (miopen_launch_status && profile_result.is_valid()) {
results.emplace_back();

14
tensorflow/core/kernels/conv_ops_3d.cc
View File

@ -377,12 +377,15 @@ struct LaunchConvOp<GPUDevice, T, OpKernelContext> {
}
}
#elif TENSORFLOW_USE_ROCM
DnnScratchAllocator scratch_allocator(ConvolveScratchSize, ctx);
std::vector<ProfileResult> algorithms;
OP_REQUIRES(ctx,
stream->parent()->GetMIOpenConvolveAlgorithms(
se::dnn::ConvolutionKind::FORWARD, stream,
se::dnn::ToDataType<T>::value, input_desc, filter_desc,
conv_desc, output_desc, &algorithms),
se::dnn::ConvolutionKind::FORWARD,
se::dnn::ToDataType<T>::value, stream, input_desc,
input_ptr, filter_desc, filter_ptr, output_desc,
output_ptr, conv_desc, &scratch_allocator, &algorithms),
errors::Unknown(
"Failed to get convolution algorithm. This is probably "
"because MIOpen failed to initialize, so try looking to "
@ -403,14 +406,15 @@ struct LaunchConvOp<GPUDevice, T, OpKernelContext> {
} else {
for (auto miopen_algorithm : algorithms) {
auto profile_algorithm = miopen_algorithm.algorithm();
DnnScratchAllocator scratch_allocator(ConvolveScratchSize, ctx);
ProfileResult profile_result;
bool miopen_launch_status =
stream
->ThenConvolveWithAlgorithm(
input_desc, input_ptr, filter_desc, filter_ptr, conv_desc,
output_desc, &output_ptr, &scratch_allocator,
AlgorithmConfig(profile_algorithm), &profile_result)
AlgorithmConfig(profile_algorithm,
miopen_algorithm.scratch_size()),
&profile_result)
.ok();
if (miopen_launch_status) {
if (profile_result.is_valid()) {

1
tensorflow/core/kernels/gpu_utils.cc
View File

@ -270,6 +270,7 @@ Status BestCudnnConvAlgorithm(absl::Span<const AutotuneResult> results,
algo->set_algorithm({selected_result->conv().algorithm(),
selected_result->conv().tensor_ops_enabled()});
algo->set_scratch_size(selected_result->scratch_bytes());
if (selected_result_no_scratch != filtered_results_no_scratch.end()) {
algo->set_algorithm_no_scratch(
{selected_result_no_scratch->conv().algorithm(),

8
tensorflow/python/debug/lib/check_numerics_callback_test.py
View File

@ -41,6 +41,7 @@ from tensorflow.python.ops import gradient_checker_v2
from tensorflow.python.ops import math_ops
from tensorflow.python.ops import variables
from tensorflow.python.platform import googletest
from tensorflow.python.platform import test as test_lib
class LimitStringLengthTest(test_util.TensorFlowTestCase):
@ -155,6 +156,13 @@ class CheckNumericsCallbackTest(test_util.TensorFlowTestCase):
@test_util.run_in_graph_and_eager_modes
def testMobileNetV2Fit(self):
"""Test training Keras MobileNetV2 application works w/ check numerics."""
if test_lib.is_built_with_rocm():
# This test passes with MIOpen Find Mode (which is the default)
# This bug is being tracked via MLOpen Issue #2379, re-enable this
# test once the fix for that issue is available in a ROCm release
self.skipTest("MIOpen bug results in test failure")
check_numerics_callback.enable_check_numerics()
model = mobilenet_v2.MobileNetV2(alpha=0.1, weights=None)

11
tensorflow/stream_executor/dnn.cc
View File

@ -48,12 +48,15 @@ bool DnnSupport::GetConvolveAlgorithms(
}
bool DnnSupport::GetMIOpenConvolveAlgorithms(
dnn::ConvolutionKind /*kind*/, Stream* /*stream*/,
dnn::DataType /*element_type*/,
const dnn::BatchDescriptor& /*input_descriptor*/,
dnn::ConvolutionKind /*kind*/, dnn::DataType /*element_type*/,
Stream* /*stream*/, const dnn::BatchDescriptor& /*input_descriptor*/,
DeviceMemoryBase input_data,
const dnn::FilterDescriptor& /*filter_descriptor*/,
const dnn::ConvolutionDescriptor& /*convolution_descriptor*/,
DeviceMemoryBase filter_data,
const dnn::BatchDescriptor& /*output_descriptor*/,
DeviceMemoryBase output_data,
const dnn::ConvolutionDescriptor& /*convolution_descriptor*/,
ScratchAllocator* scratch_allocator,
std::vector<ProfileResult>* /*out_algorithms*/) {
return false;
}

9
tensorflow/stream_executor/dnn.h
View File

@ -1346,11 +1346,14 @@ class DnnSupport {
std::vector<AlgorithmDesc>* out_algorithms);
virtual bool GetMIOpenConvolveAlgorithms(
dnn::ConvolutionKind kind, Stream* stream, dnn::DataType element_type,
const dnn::BatchDescriptor& input_descriptor,
dnn::ConvolutionKind kind, dnn::DataType element_type, Stream* stream,
const dnn::BatchDescriptor& input_descriptor, DeviceMemoryBase input_data,
const dnn::FilterDescriptor& filter_descriptor,
const dnn::ConvolutionDescriptor& convolution_descriptor,
DeviceMemoryBase filter_data,
const dnn::BatchDescriptor& output_descriptor,
DeviceMemoryBase output_data,
const dnn::ConvolutionDescriptor& convolution_descriptor,
ScratchAllocator* scratch_allocator,
std::vector<ProfileResult>* out_algorithms);
// Returns a list of supported rnn algorithms.

702
tensorflow/stream_executor/rocm/rocm_dnn.cc
View File

@ -54,12 +54,13 @@ NarrowT CheckedNarrowing(const WideT& wide) {
return narrow;
}
const int kImmediateModeVlogLevel = 3;
const int kConvDebugVlogLevel = 3;
} // namespace
namespace stream_executor {
using dnn::AlgorithmDesc;
using dnn::BatchDescriptor;
using dnn::ConvolutionDescriptor;
using dnn::FilterDescriptor;
@ -94,6 +95,75 @@ string ToString(miopenStatus_t status) {
}
}
string ToString(miopenConvFwdAlgorithm_t algorithm) {
string s;
switch (algorithm) {
case miopenConvolutionFwdAlgoGEMM:
s = "GEMM";
break;
case miopenConvolutionFwdAlgoDirect:
s = "Direct";
break;
case miopenConvolutionFwdAlgoFFT:
s = "FFT";
break;
case miopenConvolutionFwdAlgoWinograd:
s = "Winograd";
break;
case miopenConvolutionFwdAlgoImplicitGEMM:
s = "Implicit GEMM";
break;
case miopenConvolutionFwdAlgoStaticCompiledGEMM:
s = "Static Compiled GEMM";
break;
}
return s;
}
string ToString(miopenConvBwdWeightsAlgorithm_t algorithm) {
string s;
switch (algorithm) {
case miopenConvolutionBwdWeightsAlgoGEMM:
s = "GEMM";
break;
case miopenConvolutionBwdWeightsAlgoDirect:
s = "Direct";
break;
case miopenConvolutionBwdWeightsAlgoWinograd:
s = "Winograd";
break;
case miopenConvolutionBwdWeightsAlgoImplicitGEMM:
s = "Implicit GEMM";
break;
}
return s;
}
string ToString(miopenConvBwdDataAlgorithm_t algorithm) {
string s;
switch (algorithm) {
case miopenConvolutionBwdDataAlgoGEMM:
s = "GEMM";
break;
case miopenConvolutionBwdDataAlgoDirect:
s = "Direct";
break;
case miopenConvolutionBwdDataAlgoFFT:
s = "FFT";
break;
case miopenConvolutionBwdDataAlgoWinograd:
s = "Winograd";
break;
case miopenTransposeBwdDataAlgoGEMM:
s = "Transpose GEMM";
break;
case miopenConvolutionBwdDataAlgoImplicitGEMM:
s = "Implicit GEMM";
break;
}
return s;
}
string ToString(miopenConvAlgorithm_t algorithm) {
string s;
switch (algorithm) {
@ -109,9 +179,16 @@ string ToString(miopenConvAlgorithm_t algorithm) {
case miopenConvolutionAlgoWinograd:
s = "Winograd";
break;
case miopenConvolutionAlgoImplicitGEMM:
s = "Implicit GEMM";
break;
case miopenConvolutionAlgoStaticCompiledGEMM:
s = "Static Compiled GEMM";
break;
}
return s;
}
// RAII wrapper for all calls to MIOpen with a MIOpen handle argument.
//
// See MIOpenAccess::GetHandle() for details.
@ -433,66 +510,38 @@ std::set<uint64> CachedFusionPlans::unsupported_plans;
dnn::ProfileResult GetProfileResultFromConvSolution(
miopenConvSolution_t solution) {
dnn::ProfileResult profile_result;
profile_result.set_algorithm({solution.solution_id, false});
profile_result.set_algorithm(
{static_cast<AlgorithmDesc::Index>(solution.solution_id), false});
profile_result.set_elapsed_time_in_ms(solution.time);
profile_result.set_scratch_size(solution.workspace_size);
return profile_result;
}
} // namespace
namespace {
miopenHandle_t ToHandle(void* opaque_handle) {
return static_cast<miopenHandle_t>(opaque_handle);
}
miopenConvFwdAlgorithm_t ToConvForwardAlgo(dnn::AlgorithmDesc algorithm) {
miopenConvFwdAlgorithm_t algo = miopenConvFwdAlgorithm_t(algorithm.algo_id());
switch (algo) {
case miopenConvolutionFwdAlgoGEMM:
case miopenConvolutionFwdAlgoDirect:
case miopenConvolutionFwdAlgoFFT:
case miopenConvolutionFwdAlgoWinograd:
return algo;
dnn::ProfileResult GetProfileResultFromConvAlgoPerf(
dnn::ConvolutionKind kind, miopenConvAlgoPerf_t algorithm) {
dnn::ProfileResult profile_result;
switch (kind) {
case dnn::ConvolutionKind::FORWARD:
profile_result.set_algorithm(
{static_cast<AlgorithmDesc::Index>(algorithm.fwd_algo), false});
break;
case dnn::ConvolutionKind::BACKWARD_DATA:
profile_result.set_algorithm(
{static_cast<AlgorithmDesc::Index>(algorithm.bwd_data_algo), false});
break;
case dnn::ConvolutionKind::BACKWARD_FILTER:
profile_result.set_algorithm(
{static_cast<AlgorithmDesc::Index>(algorithm.bwd_weights_algo),
false});
break;
default:
LOG(FATAL) << "Unsupported MIOpen convolution forward algorithm: "
<< algorithm.algo_id();
LOG(FATAL) << "Unexpected convolution kind " << static_cast<int>(kind);
break;
}
profile_result.set_elapsed_time_in_ms(algorithm.time);
profile_result.set_scratch_size(algorithm.memory);
return profile_result;
}
miopenConvBwdDataAlgorithm_t ToConvBackwardDataAlgo(
dnn::AlgorithmDesc algorithm) {
miopenConvBwdDataAlgorithm_t algo =
miopenConvBwdDataAlgorithm_t(algorithm.algo_id());
switch (algo) {
case miopenConvolutionBwdDataAlgoGEMM:
case miopenConvolutionBwdDataAlgoDirect:
case miopenConvolutionBwdDataAlgoFFT:
case miopenConvolutionBwdDataAlgoWinograd:
return algo;
default:
LOG(FATAL)
<< "Unsupported MIOpen convolution backward algorithm for data: "
<< algorithm.algo_id();
}
}
miopenConvBwdWeightsAlgorithm_t ToConvBackwardFilterAlgo(
dnn::AlgorithmDesc algorithm) {
miopenConvBwdWeightsAlgorithm_t algo =
miopenConvBwdWeightsAlgorithm_t(algorithm.algo_id());
switch (algo) {
case miopenConvolutionBwdWeightsAlgoGEMM:
case miopenConvolutionBwdWeightsAlgoDirect:
return algo;
default:
LOG(FATAL)
<< "Unsupported MIOpen convolution backward algorithm for filter: "
<< algorithm.algo_id();
}
}
} // namespace
// Wraps a MIOpen handle and provides access to it through miopenHandle_t
@ -542,7 +591,21 @@ class MIOpenAccess {
miopenHandle_t handle_ TF_GUARDED_BY(mutex_); // Owned.
};
MIOpenSupport::MIOpenSupport(GpuExecutor* parent) : parent_(parent) {}
MIOpenSupport::MIOpenSupport(GpuExecutor* parent) : parent_(parent) {
// by default, the Get*Algorithm API will return the list of all applicable
// algorithms
return_best_algo_only_ = false;
// but if the env var TF_ROCM_RETURN_BEST_ALGO_ONLY is set, only the best
// (i.e. most efficient) algorithm will be returned
tensorflow::ReadBoolFromEnvVar("TF_ROCM_RETURN_BEST_ALGO_ONLY", false,
&return_best_algo_only_);
// by default, use Find Mode APIs for convolution
use_immediate_mode_ = false;
// swich to Find Mode if env var TF_ROCM_USE_IMMEDIATE_MODE is set
tensorflow::ReadBoolFromEnvVar("TF_ROCM_USE_IMMEDIATE_MODE", false,
&use_immediate_mode_);
}
port::Status MIOpenSupport::Init() {
ScopedActivateExecutorContext context(parent_);
@ -1593,11 +1656,6 @@ miopenDataType_t ToMIOpenDataType(
}
}
miopenDataType_t ToMIOpenDataType(dnn::DataType data_type,
dnn::FilterLayout filter_layout) {
return ToMIOpenDataType(data_type);
}
miopenRNNInputMode_t ToMIOpenRnnInputMode(dnn::RnnInputMode input_mode) {
switch (input_mode) {
case dnn::RnnInputMode::kRnnLinearSkip:
@ -1637,38 +1695,11 @@ miopenRNNMode_t ToMIOpenRnnMode(dnn::RnnMode rnn_mode) {
}
}
int MIOpenDataTypeToByteSize(miopenDataType_t data_type) {
switch (data_type) {
case miopenFloat:
return sizeof(float);
case miopenHalf:
return sizeof(Eigen::half);
default:
LOG(FATAL) << "Invalid DNN data type: " << static_cast<int>(data_type);
}
}
template <typename Base>
class MixinBase : public Base {};
template <>
class MixinBase<void> {};
dnn::DataType GetConvAccumulatorType(dnn::DataType data_type) {
switch (data_type) {
case dnn::DataType::kFloat:
case dnn::DataType::kDouble:
return data_type;
case dnn::DataType::kHalf:
// FIXME: Check if MIOpen can switch dynamically change accumulator type
return dnn::DataType::kFloat;
case dnn::DataType::kInt8:
case dnn::DataType::kInt32:
return dnn::DataType::kInt32;
default:
LOG(FATAL) << "Invalid DNN data type: " << static_cast<int>(data_type);
}
}
} // namespace
#define RETURN_IF_MIOPEN_ERROR(STATUS, ...) \
@ -2818,21 +2849,6 @@ port::Status MIOpenSupport::DoPrepareForConvolution(
const dnn::AlgorithmConfig& algorithm_config,
ScratchAllocator* scratch_allocator, dnn::AlgorithmDesc* algorithm_desc,
DeviceMemory<uint8>* scratch_memory) {
ScopedTensorDescriptor input_nd{
input_descriptor,
ToMIOpenDataType(element_type, input_descriptor.layout())};
ScopedFilterDescriptor filter{
filter_descriptor, input_descriptor,
ToMIOpenDataType(element_type, filter_descriptor.layout())};
ScopedTensorDescriptor output_nd{
output_descriptor,
ToMIOpenDataType(element_type, output_descriptor.layout())};
ScopedConvolutionDescriptor conv{
convolution_descriptor,
ToMIOpenDataType(GetConvAccumulatorType(element_type))};
auto miopen = miopen_->GetHandle(parent_, stream);
absl::optional<dnn::AlgorithmDesc> input_algo_desc =
algorithm_config.algorithm();
@ -2841,64 +2857,9 @@ port::Status MIOpenSupport::DoPrepareForConvolution(
// An algorithm has been specified.
*algorithm_desc = *input_algo_desc;
const uint64_t solution_id = algorithm_desc->algo_id();
assert(algorithm_config.scratch_size().has_value());
size_t scratch_memory_size = 0;
switch (kind) {
case dnn::ConvolutionKind::FORWARD: {
auto status = wrap::miopenConvolutionForwardGetSolutionWorkspaceSize(
miopen.handle(), filter.handle(), input_nd.handle(), conv.handle(),
output_nd.handle(), solution_id, &scratch_memory_size);
if (status != miopenStatusSuccess) {
return port::InternalError(absl::StrCat(
"call to miopenConvolutionForwardGetSolutionWorkspaceSize "
"failed: ",
ToString(status)));
}
break;
}
case dnn::ConvolutionKind::BACKWARD_DATA: {
auto status = wrap::miopenConvolutionBackwardDataGetSolutionWorkspaceSize(
miopen.handle(), output_nd.handle(), filter.handle(), conv.handle(),
input_nd.handle(), solution_id, &scratch_memory_size);
if (status != miopenStatusSuccess) {
return port::InternalError(absl::StrCat(
"call to miopenConvolutionabckwardDataGetSolutionWorkspaceSize "
"failed: ",
ToString(status)));
}
break;
}
case dnn::ConvolutionKind::BACKWARD_FILTER: {
auto status =
wrap::miopenConvolutionBackwardWeightsGetSolutionWorkspaceSize(
miopen.handle(), output_nd.handle(), input_nd.handle(),
conv.handle(), filter.handle(), solution_id,
&scratch_memory_size);
if (status != miopenStatusSuccess) {
return port::InternalError(absl::StrCat(
"call to miopenConvolutionabckwardWeightsGetSolutionWorkspaceSize "
"failed: ",
ToString(status)));
}
break;
}
default: {
return port::InternalError(
absl::StrCat("Unexpected convolution kind ", static_cast<int>(kind)));
break;
}
}
VLOG(2) << "miopen...GetSolutionWorkspaceSize returned "
<< scratch_memory_size << " for solution_id " << solution_id;
size_t scratch_memory_size = *(algorithm_config.scratch_size());
// allocate scratch memory
if (scratch_memory_size != 0) {
@ -3015,16 +2976,24 @@ port::Status MIOpenSupport::DoConvolve(
}
}
const uint64_t solution_id = algorithm_desc.algo_id();
miopenStatus_t status = miopenStatusSuccess;
switch (kind) {
case dnn::ConvolutionKind::FORWARD: {
status = wrap::miopenConvolutionForwardImmediate(
miopen.handle(), filter.handle(), filter_data.opaque(),
input_nd.handle(), input_data.opaque(), conv.handle(),
output_nd.handle(), output_data.opaque(), scratch_memory.opaque(),
scratch_memory.size(), solution_id);
if (use_immediate_mode_) {
status = wrap::miopenConvolutionForwardImmediate(
miopen.handle(), filter.handle(), filter_data.opaque(),
input_nd.handle(), input_data.opaque(), conv.handle(),
output_nd.handle(), output_data.opaque(), scratch_memory.opaque(),
scratch_memory.size(),
static_cast<uint64_t>(algorithm_desc.algo_id()));
} else {
status = wrap::miopenConvolutionForward(
miopen.handle(), &alpha, input_nd.handle(), input_data.opaque(),
filter.handle(), filter_data.opaque(), conv.handle(),
static_cast<miopenConvFwdAlgorithm_t>(algorithm_desc.algo_id()),
&beta, output_nd.handle(), output_data.opaque(),
scratch_memory.opaque(), scratch_memory.size());
}
break;
}
@ -3037,11 +3006,21 @@ port::Status MIOpenSupport::DoConvolve(
stream, miopen.handle(), ToMIOpenDataType(element_type),
&output_back_descriptor, output_data, &transform_scratch);
status = wrap::miopenConvolutionBackwardDataImmediate(
miopen.handle(), output_nd.handle(), output_data.opaque(),
filter.handle(), filter_data.opaque(), conv.handle(),
input_nd.handle(), input_data.opaque(), scratch_memory.opaque(),
scratch_memory.size(), solution_id);
if (use_immediate_mode_) {
status = wrap::miopenConvolutionBackwardDataImmediate(
miopen.handle(), output_nd.handle(), output_data.opaque(),
filter.handle(), filter_data.opaque(), conv.handle(),
input_nd.handle(), input_data.opaque(), scratch_memory.opaque(),
scratch_memory.size(),
static_cast<uint64_t>(algorithm_desc.algo_id()));
} else {
status = wrap::miopenConvolutionBackwardData(
miopen.handle(), &alpha, output_nd.handle(), output_data.opaque(),
filter.handle(), filter_data.opaque(), conv.handle(),
static_cast<miopenConvBwdDataAlgorithm_t>(algorithm_desc.algo_id()),
&beta, input_nd.handle(), input_data.opaque(),
scratch_memory.opaque(), scratch_memory.size());
}
break;
}
case dnn::ConvolutionKind::BACKWARD_FILTER: {
@ -3053,11 +3032,22 @@ port::Status MIOpenSupport::DoConvolve(
stream, miopen.handle(), ToMIOpenDataType(element_type),
&output_back_descriptor, output_data, &transform_scratch);
status = wrap::miopenConvolutionBackwardWeightsImmediate(
miopen.handle(), output_nd.handle(), output_data.opaque(),
input_nd.handle(), input_data.opaque(), conv.handle(),
filter.handle(), filter_data.opaque(), scratch_memory.opaque(),
scratch_memory.size(), solution_id);
if (use_immediate_mode_) {
status = wrap::miopenConvolutionBackwardWeightsImmediate(
miopen.handle(), output_nd.handle(), output_data.opaque(),
input_nd.handle(), input_data.opaque(), conv.handle(),
filter.handle(), filter_data.opaque(), scratch_memory.opaque(),
scratch_memory.size(),
static_cast<uint64_t>(algorithm_desc.algo_id()));
} else {
status = wrap::miopenConvolutionBackwardWeights(
miopen.handle(), &alpha, output_nd.handle(), output_data.opaque(),
input_nd.handle(), input_data.opaque(), conv.handle(),
static_cast<miopenConvBwdWeightsAlgorithm_t>(
algorithm_desc.algo_id()),
&beta, filter.handle(), filter_data.opaque(),
scratch_memory.opaque(), scratch_memory.size());
}
break;
}
default:
@ -3104,11 +3094,35 @@ bool MIOpenSupport::GetConvolveAlgorithms(
}
bool MIOpenSupport::GetMIOpenConvolveAlgorithms(
dnn::ConvolutionKind kind, Stream* stream, dnn::DataType element_type,
const dnn::BatchDescriptor& input_descriptor,
dnn::ConvolutionKind kind, dnn::DataType element_type, Stream* stream,
const dnn::BatchDescriptor& input_descriptor, DeviceMemoryBase input_data,
const dnn::FilterDescriptor& filter_descriptor,
DeviceMemoryBase filter_data, const dnn::BatchDescriptor& output_descriptor,
DeviceMemoryBase output_data,
const dnn::ConvolutionDescriptor& convolution_descriptor,
const dnn::BatchDescriptor& output_descriptor,
ScratchAllocator* scratch_allocator,
std::vector<dnn::ProfileResult>* out_algorithms) {
return use_immediate_mode_
? GetMIOpenConvolveAlgorithmsImmediateMode(
kind, element_type, stream, input_descriptor, input_data,
filter_descriptor, filter_data, output_descriptor,
output_data, convolution_descriptor, scratch_allocator,
out_algorithms)
: GetMIOpenConvolveAlgorithmsFindMode(
kind, element_type, stream, input_descriptor, input_data,
filter_descriptor, filter_data, output_descriptor,
output_data, convolution_descriptor, scratch_allocator,
out_algorithms);
}
bool MIOpenSupport::GetMIOpenConvolveAlgorithmsImmediateMode(
dnn::ConvolutionKind kind, dnn::DataType element_type, Stream* stream,
const dnn::BatchDescriptor& input_descriptor, DeviceMemoryBase input_data,
const dnn::FilterDescriptor& filter_descriptor,
DeviceMemoryBase filter_data, const dnn::BatchDescriptor& output_descriptor,
DeviceMemoryBase output_data,
const dnn::ConvolutionDescriptor& convolution_descriptor,
ScratchAllocator* scratch_allocator,
std::vector<dnn::ProfileResult>* out_algorithms) {
auto miopen = miopen_->GetHandle(parent_, stream);
@ -3169,18 +3183,14 @@ bool MIOpenSupport::GetMIOpenConvolveAlgorithms(
}
}
VLOG(kImmediateModeVlogLevel)
VLOG(kConvDebugVlogLevel)
<< "Number of conv solutions max: " << maxSolutionCount;
// if the env var TF_ROCM_MIMIC_FIND_MODE is set, determine the best solution
// as per the "runtime" information for each solution (returned by the prior
// call to the *GetSolution api), and then return only the best solution
// The idea here is to mimic the old "find" mode, in which we relied upon
// the miopen api to determine the best solution, and use that solution
// without doing any further measurement in the TF layer
bool mimic_find_mode = false;
tensorflow::ReadBoolFromEnvVar("TF_ROCM_MIMIC_FIND_MODE", false,
&mimic_find_mode);
if (return_best_algo_only_) {
VLOG(kConvDebugVlogLevel) << "TF_ROCM_RETURN_BEST_ALGO_ONLY is set, "
<< "setting maxSolutionCount to 1";
maxSolutionCount = 1;
}
size_t solutionCount = 0;
std::unique_ptr<miopenConvSolution_t[]> solutions(
@ -3199,61 +3209,30 @@ bool MIOpenSupport::GetMIOpenConvolveAlgorithms(
return false;
}
VLOG(kImmediateModeVlogLevel)
VLOG(kConvDebugVlogLevel)
<< "Number of conv solutions actual: " << solutionCount;
if (mimic_find_mode) {
miopenConvSolution_t best_solution = solutions[0];
for (size_t i = 0; i < solutionCount; i++) {
miopenConvSolution_t solution = solutions[i];
for (int i = 1; i < solutionCount; i++) {
miopenConvSolution_t solution = solutions[i];
if (solution.time < best_solution.time) {
best_solution = solution;
}
}
VLOG(kImmediateModeVlogLevel)
<< "Best Solution (id, algo) = " << best_solution.solution_id
<< ", " << ToString(best_solution.algorithm);
VLOG(kConvDebugVlogLevel)
<< "solution " << i << " (time, mem, id, algo) = " << solution.time
<< ", " << solution.workspace_size << ", " << solution.solution_id
<< ", " << ToString(solution.algorithm);
status = wrap::miopenConvolutionForwardCompileSolution(
miopen.handle(), filter.handle(), input_nd.handle(), conv.handle(),
output_nd.handle(), best_solution.solution_id);
output_nd.handle(), solution.solution_id);
if (status != miopenStatusSuccess) {
LOG(FATAL) << "call to miopenConvolutionForwardCompileSolution "
"failed: "
<< ToString(status);
LOG(FATAL)
<< "call to miopenConvolutionForwardCompileSolution failed: "
<< ToString(status);
return false;
}
out_algorithms->emplace_back(
GetProfileResultFromConvSolution(best_solution));
} else {
for (int i = 0; i < solutionCount; i++) {
miopenConvSolution_t solution = solutions[i];
VLOG(kImmediateModeVlogLevel)
<< "solution " << i
<< " (time, mem, id, algo) = " << solution.time << ", "
<< solution.workspace_size << ", " << solution.solution_id << ", "
<< ToString(solution.algorithm);
status = wrap::miopenConvolutionForwardCompileSolution(
miopen.handle(), filter.handle(), input_nd.handle(),
conv.handle(), output_nd.handle(), solution.solution_id);
if (status != miopenStatusSuccess) {
LOG(FATAL)
<< "call to miopenConvolutionForwardCompileSolution failed: "
<< ToString(status);
return false;
}
out_algorithms->emplace_back(
GetProfileResultFromConvSolution(solution));
}
GetProfileResultFromConvSolution(solution));
}
break;
}
@ -3269,62 +3248,30 @@ bool MIOpenSupport::GetMIOpenConvolveAlgorithms(
return false;
}
VLOG(kImmediateModeVlogLevel)
VLOG(kConvDebugVlogLevel)
<< "Number of conv solutions actual: " << solutionCount;
if (mimic_find_mode) {
miopenConvSolution_t best_solution = solutions[0];
for (size_t i = 0; i < solutionCount; i++) {
miopenConvSolution_t solution = solutions[i];
for (int i = 1; i < solutionCount; i++) {
miopenConvSolution_t solution = solutions[i];
if (solution.time < best_solution.time) {
best_solution = solution;
}
}
VLOG(kImmediateModeVlogLevel)
<< "Best Solution (id, algo) = " << best_solution.solution_id
<< ", " << ToString(best_solution.algorithm);
VLOG(kConvDebugVlogLevel)
<< "solution " << i << " (time, mem, id, algo) = " << solution.time
<< ", " << solution.workspace_size << ", " << solution.solution_id
<< ", " << ToString(solution.algorithm);
status = wrap::miopenConvolutionBackwardDataCompileSolution(
miopen.handle(), output_nd.handle(), filter.handle(), conv.handle(),
input_nd.handle(), best_solution.solution_id);
input_nd.handle(), solution.solution_id);
if (status != miopenStatusSuccess) {
LOG(FATAL) << "call to miopenConvolutionBackwardDataCompileSolution "
LOG(FATAL) << " call to miopenConvolutionBackwardDataCompileSolution "
"failed: "
<< ToString(status);
return false;
}
out_algorithms->emplace_back(
GetProfileResultFromConvSolution(best_solution));
} else {
for (int i = 0; i < solutionCount; i++) {
miopenConvSolution_t solution = solutions[i];
VLOG(kImmediateModeVlogLevel)
<< "solution " << i
<< " (time, mem, id, algo) = " << solution.time << ", "
<< solution.workspace_size << ", " << solution.solution_id << ", "
<< ToString(solution.algorithm);
status = wrap::miopenConvolutionBackwardDataCompileSolution(
miopen.handle(), output_nd.handle(), filter.handle(),
conv.handle(), input_nd.handle(), solution.solution_id);
if (status != miopenStatusSuccess) {
LOG(FATAL)
<< " call to miopenConvolutionBackwardDataCompileSolution "
"failed: "
<< ToString(status);
return false;
}
out_algorithms->emplace_back(
GetProfileResultFromConvSolution(solution));
}
GetProfileResultFromConvSolution(solution));
}
break;
}
@ -3339,26 +3286,20 @@ bool MIOpenSupport::GetMIOpenConvolveAlgorithms(
return false;
}
VLOG(kImmediateModeVlogLevel)
VLOG(kConvDebugVlogLevel)
<< "Number of conv solutions actual: " << solutionCount;
if (mimic_find_mode) {
miopenConvSolution_t best_solution = solutions[0];
for (size_t i = 0; i < solutionCount; i++) {
miopenConvSolution_t solution = solutions[i];
for (int i = 1; i < solutionCount; i++) {
miopenConvSolution_t solution = solutions[i];
if (solution.time < best_solution.time) {
best_solution = solution;
}
}
VLOG(kImmediateModeVlogLevel)
<< "Best Solution (id, algo) = " << best_solution.solution_id
<< ", " << ToString(best_solution.algorithm);
VLOG(kConvDebugVlogLevel)
<< "solution " << i << " (time, mem, id, algo) = " << solution.time
<< ", " << solution.workspace_size << ", " << solution.solution_id
<< ", " << ToString(solution.algorithm);
status = wrap::miopenConvolutionBackwardWeightsCompileSolution(
miopen.handle(), output_nd.handle(), input_nd.handle(),
conv.handle(), filter.handle(), best_solution.solution_id);
conv.handle(), filter.handle(), solution.solution_id);
if (status != miopenStatusSuccess) {
LOG(FATAL)
@ -3369,33 +3310,7 @@ bool MIOpenSupport::GetMIOpenConvolveAlgorithms(
}
out_algorithms->emplace_back(
GetProfileResultFromConvSolution(best_solution));
} else {
for (int i = 0; i < solutionCount; i++) {
miopenConvSolution_t solution = solutions[i];
VLOG(kImmediateModeVlogLevel)
<< "solution " << i
<< " (time, mem, id, algo) = " << solution.time << ", "
<< solution.workspace_size << ", " << solution.solution_id << ", "
<< ToString(solution.algorithm);
status = wrap::miopenConvolutionBackwardWeightsCompileSolution(
miopen.handle(), output_nd.handle(), input_nd.handle(),
conv.handle(), filter.handle(), solution.solution_id);
if (status != miopenStatusSuccess) {
LOG(FATAL)
<< "call to miopenConvolutionBackwardWeightsCompileSolution "
"failed: "
<< ToString(status);
return false;
}
out_algorithms->emplace_back(
GetProfileResultFromConvSolution(solution));
}
GetProfileResultFromConvSolution(solution));
}
break;
}
@ -3409,6 +3324,165 @@ bool MIOpenSupport::GetMIOpenConvolveAlgorithms(
return true;
}
bool MIOpenSupport::GetMIOpenConvolveAlgorithmsFindMode(
dnn::ConvolutionKind kind, dnn::DataType element_type, Stream* stream,
const dnn::BatchDescriptor& input_descriptor, DeviceMemoryBase input_data,
const dnn::FilterDescriptor& filter_descriptor,
DeviceMemoryBase filter_data, const dnn::BatchDescriptor& output_descriptor,
DeviceMemoryBase output_data,
const dnn::ConvolutionDescriptor& convolution_descriptor,
ScratchAllocator* scratch_allocator,
std::vector<dnn::ProfileResult>* out_algorithms) {
auto miopen = miopen_->GetHandle(parent_, stream);
ScopedTensorDescriptor input_nd{input_descriptor,
ToMIOpenDataType(element_type)};
ScopedTensorDescriptor output_nd{output_descriptor,
ToMIOpenDataType(element_type)};
ScopedFilterDescriptor filter{filter_descriptor, input_descriptor,
ToMIOpenDataType(element_type)};
ScopedConvolutionDescriptor conv{convolution_descriptor,
ToMIOpenDataType(element_type)};
// Determine the workspace memory size that will need by the call to Find
size_t scratch_memory_size = 0;
switch (kind) {
case dnn::ConvolutionKind::FORWARD: {
auto status = wrap::miopenConvolutionForwardGetWorkSpaceSize(
miopen.handle(), filter.handle(), input_nd.handle(), conv.handle(),
output_nd.handle(), &scratch_memory_size);
if (status != miopenStatusSuccess) {
LOG(FATAL)
<< "call to miopenConvolutionForwardGetWorkspaceSize failed: "
<< ToString(status);
return false;
}
break;
}
case dnn::ConvolutionKind::BACKWARD_DATA: {
auto status = wrap::miopenConvolutionBackwardDataGetWorkSpaceSize(
miopen.handle(), output_nd.handle(), filter.handle(), conv.handle(),
input_nd.handle(), &scratch_memory_size);
if (status != miopenStatusSuccess) {
LOG(FATAL)
<< "call to miopenConvolutionBackwardDataGetWorkspaceSize failed: "
<< ToString(status);
return false;
}
break;
}
case dnn::ConvolutionKind::BACKWARD_FILTER: {
auto status = wrap::miopenConvolutionBackwardWeightsGetWorkSpaceSize(
miopen.handle(), output_nd.handle(), input_nd.handle(), conv.handle(),
filter.handle(), &scratch_memory_size);
if (status != miopenStatusSuccess) {
LOG(FATAL)
<< "call to miopenConvolutionBackwardWeightsGetWorkspaceSize "
"failed: "
<< ToString(status);
return false;
}
break;
}
default: {
LOG(FATAL) << "Unexpected convolution kind " << static_cast<int>(kind);
return false;
break;
}
}
// allocate scratch memory
DeviceMemory<uint8> scratch_memory;
if (scratch_memory_size != 0) {
if (scratch_allocator == nullptr) {
LOG(FATAL)
<< "An allocator must be specified when scratch memory is needed";
return false;
}
auto allocated = scratch_allocator->AllocateBytes(scratch_memory_size);
if (allocated.ok()) {
scratch_memory = allocated.ValueOrDie();
} else {
LOG(FATAL)
<< "Failed to allocate scratch memory - "
<< allocated.status().error_message() << "\n"
<< "\tYou can set the env var TF_CUDNN_WORKSPACE_LIMIT_IN_MB to a "
"larger number (e.g. 8192) to increase the max memory limit.\n"
<< "\tIncreasing the max memory limit might help resolve this "
"error";
return false;
}
}
// Only get the best algorithm for Find Mode
size_t requestedAlgorithmCount = 1;
VLOG(kConvDebugVlogLevel)
<< "Number of conv algortihms to request: " << requestedAlgorithmCount;
miopenConvAlgoPerf_t returnedAlgorithm;
int returnedAlgorithmCount = 0;
bool exhaustiveSearch = false;
switch (kind) {
case dnn::ConvolutionKind::FORWARD: {
auto status = wrap::miopenFindConvolutionForwardAlgorithm(
miopen.handle(), input_nd.handle(), input_data.opaque(),
filter.handle(), filter_data.opaque(), conv.handle(),
output_nd.handle(), output_data.opaque(), requestedAlgorithmCount,
&returnedAlgorithmCount, &returnedAlgorithm, scratch_memory.opaque(),
scratch_memory_size, exhaustiveSearch);
if (status != miopenStatusSuccess) {
LOG(FATAL) << "call to miopenFindConvolutionForwardAlgorithm failed: "
<< ToString(status);
return false;
}
break;
}
case dnn::ConvolutionKind::BACKWARD_DATA: {
auto status = wrap::miopenFindConvolutionBackwardDataAlgorithm(
miopen.handle(), output_nd.handle(), output_data.opaque(),
filter.handle(), filter_data.opaque(), conv.handle(),
input_nd.handle(), input_data.opaque(), requestedAlgorithmCount,
&returnedAlgorithmCount, &returnedAlgorithm, scratch_memory.opaque(),
scratch_memory_size, exhaustiveSearch);
if (status != miopenStatusSuccess) {
LOG(FATAL)
<< "call to miopenFindConvolutionBackwardDataAlgorithm failed: "
<< ToString(status);
return false;
}
break;
}
case dnn::ConvolutionKind::BACKWARD_FILTER: {
auto status = wrap::miopenFindConvolutionBackwardWeightsAlgorithm(
miopen.handle(), output_nd.handle(), output_data.opaque(),
input_nd.handle(), input_data.opaque(), conv.handle(),
filter.handle(), filter_data.opaque(), requestedAlgorithmCount,
&returnedAlgorithmCount, &returnedAlgorithm, scratch_memory.opaque(),
scratch_memory_size, exhaustiveSearch);
if (status != miopenStatusSuccess) {
LOG(FATAL) << "call to miopenConvolutionBackwardWeightsAlgorithm "
"failed: "
<< ToString(status);
return false;
}
break;
}
default: {
LOG(FATAL) << "Unexpected convolution kind " << static_cast<int>(kind);
return false;
break;
}
}
out_algorithms->emplace_back(
GetProfileResultFromConvAlgoPerf(kind, returnedAlgorithm));
return true;
}
bool MIOpenSupport::GetRnnAlgorithms(
std::vector<dnn::AlgorithmDesc>* out_algorithms) {
// ROCM TODO: implement this with proper MIOpen API

38
tensorflow/stream_executor/rocm/rocm_dnn.h
View File

@ -198,11 +198,14 @@ class MIOpenSupport : public dnn::DnnSupport {
std::vector<dnn::AlgorithmDesc>* out_algorithms) override;
bool GetMIOpenConvolveAlgorithms(
dnn::ConvolutionKind kind, Stream* stream, dnn::DataType element_type,
const dnn::BatchDescriptor& input_descriptor,
dnn::ConvolutionKind kind, dnn::DataType element_type, Stream* stream,
const dnn::BatchDescriptor& input_descriptor, DeviceMemoryBase input_data,
const dnn::FilterDescriptor& filter_descriptor,
const dnn::ConvolutionDescriptor& convolution_descriptor,
DeviceMemoryBase filter_data,
const dnn::BatchDescriptor& output_descriptor,
DeviceMemoryBase output_data,
const dnn::ConvolutionDescriptor& convolution_descriptor,
ScratchAllocator* scratch_allocator,
std::vector<dnn::ProfileResult>* out_algorithms) override;
bool GetRnnAlgorithms(
@ -650,6 +653,13 @@ class MIOpenSupport : public dnn::DnnSupport {
private:
GpuExecutor* parent_; // Parent executor object. Not owned.
// Flag to indicate whether Get*Algorithm routines should only return
// the best algorithm (as opposed to a list of all applicable ones)
bool return_best_algo_only_;
// Flag to indicate whether to use Immediate (or Find) mode for Convolutions
bool use_immediate_mode_;
// Provide access to the MIOpen handle.
std::unique_ptr<class MIOpenAccess> miopen_;
@ -814,6 +824,28 @@ class MIOpenSupport : public dnn::DnnSupport {
ScratchAllocator* scratch_allocator,
DeviceMemory<uint8>* scratch_memory) override;
bool GetMIOpenConvolveAlgorithmsImmediateMode(
dnn::ConvolutionKind kind, dnn::DataType element_type, Stream* stream,
const dnn::BatchDescriptor& input_descriptor, DeviceMemoryBase input_data,
const dnn::FilterDescriptor& filter_descriptor,
DeviceMemoryBase filter_data,
const dnn::BatchDescriptor& output_descriptor,
DeviceMemoryBase output_data,
const dnn::ConvolutionDescriptor& convolution_descriptor,
ScratchAllocator* scratch_allocator,
std::vector<dnn::ProfileResult>* out_algorithms);
bool GetMIOpenConvolveAlgorithmsFindMode(
dnn::ConvolutionKind kind, dnn::DataType element_type, Stream* stream,
const dnn::BatchDescriptor& input_descriptor, DeviceMemoryBase input_data,
const dnn::FilterDescriptor& filter_descriptor,
DeviceMemoryBase filter_data,
const dnn::BatchDescriptor& output_descriptor,
DeviceMemoryBase output_data,
const dnn::ConvolutionDescriptor& convolution_descriptor,
ScratchAllocator* scratch_allocator,
std::vector<dnn::ProfileResult>* out_algorithms);
SE_DISALLOW_COPY_AND_ASSIGN(MIOpenSupport);
};

13
tensorflow/stream_executor/stream_executor_pimpl.cc
View File

@ -291,19 +291,22 @@ bool StreamExecutor::GetConvolveAlgorithms(
}
bool StreamExecutor::GetMIOpenConvolveAlgorithms(
dnn::ConvolutionKind kind, Stream *stream, dnn::DataType element_type,
const dnn::BatchDescriptor &input_descriptor,
dnn::ConvolutionKind kind, dnn::DataType element_type, Stream *stream,
const dnn::BatchDescriptor &input_descriptor, DeviceMemoryBase input_data,
const dnn::FilterDescriptor &filter_descriptor,
DeviceMemoryBase filter_data, const dnn::BatchDescriptor &output_descriptor,
DeviceMemoryBase output_data,
const dnn::ConvolutionDescriptor &convolution_descriptor,
const dnn::BatchDescriptor &output_descriptor,
ScratchAllocator *scratch_allocator,
std::vector<dnn::ProfileResult> *out_algorithms) {
dnn::DnnSupport *dnn_support = AsDnn();
if (!dnn_support) {
return false;
}
return dnn_support->GetMIOpenConvolveAlgorithms(
kind, stream, element_type, input_descriptor, filter_descriptor,
convolution_descriptor, output_descriptor, out_algorithms);
kind, element_type, stream, input_descriptor, input_data,
filter_descriptor, filter_data, output_descriptor, output_data,
convolution_descriptor, scratch_allocator, out_algorithms);
}
bool StreamExecutor::GetRnnAlgorithms(

9
tensorflow/stream_executor/stream_executor_pimpl.h
View File

@ -376,11 +376,14 @@ class StreamExecutor {
// Returns the list of supported algorithms for the forward convolution
// operation.
bool GetMIOpenConvolveAlgorithms(
dnn::ConvolutionKind kind, Stream *stream, dnn::DataType element_type,
const dnn::BatchDescriptor &input_descriptor,
dnn::ConvolutionKind kind, dnn::DataType element_type, Stream *stream,
const dnn::BatchDescriptor &input_descriptor, DeviceMemoryBase input_data,
const dnn::FilterDescriptor &filter_descriptor,
const dnn::ConvolutionDescriptor &convolution_descriptor,
DeviceMemoryBase filter_data,
const dnn::BatchDescriptor &output_descriptor,
DeviceMemoryBase output_data,
const dnn::ConvolutionDescriptor &convolution_descriptor,
ScratchAllocator *scratch_allocator,
std::vector<dnn::ProfileResult> *out_algorithms);
// Returns the list of supported algorithms for rnn operation.