Skip to content

Make FLAGS_determinstic effective in conv2d forward. #37173

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Nov 16, 2021
Merged

Make FLAGS_determinstic effective in conv2d forward. #37173

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
024ff37
Make FLAGS_determinstic effective in conv2d forward.
Xreki Nov 13, 2021
636f4ed
Merge branch 'develop' into fix_conv_algo
Xreki Nov 15, 2021
7270b38
Add call of SetCinnCudnnDeterministic in cinn_launch op.
Xreki Nov 15, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 8 additions & 0 deletions paddle/fluid/operators/cinn_launch_op.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,6 +15,8 @@
#include "paddle/fluid/operators/cinn_launch_op.h"
#include "paddle/fluid/string/string_helper.h"

DECLARE_bool(cudnn_deterministic);

namespace paddle {
namespace operators {

Expand Down Expand Up @@ -67,6 +69,12 @@ void LaunchCinnExecution(const CinnCompiledObject& compiled_obj,
compiled_obj.runtime_program->Execute(&context.FinalizeArguments());
}

void SetCinnRuntimeFlags() {
VLOG(4) << "Set FLAGS_cinn_cudnn_deterministic to "
<< FLAGS_cudnn_deterministic;
::cinn::runtime::SetCinnCudnnDeterministic(FLAGS_cudnn_deterministic);
}

CinnLaunchContext::CinnLaunchContext(const CinnCompiledObject& compiled_obj)
: paddle2cinn_varmap_(compiled_obj.paddle2cinn_varmap),
cinn_scope_(compiled_obj.scope) {
Expand Down
9 changes: 8 additions & 1 deletion paddle/fluid/operators/cinn_launch_op.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,7 @@
#include "cinn/hlir/framework/graph_compiler.h"
#include "cinn/hlir/framework/scope.h"
#include "cinn/runtime/cinn_runtime.h"
#include "cinn/runtime/flags.h"
#include "paddle/fluid/framework/data_type.h"
#include "paddle/fluid/framework/op_registry.h"
#include "paddle/fluid/framework/operator.h"
Expand Down Expand Up @@ -110,6 +111,9 @@ void DebugCinnCompiledResult(const CinnCompiledObject& result);
// Launch cinn to execute compiled executable program and wait done
void LaunchCinnExecution(const CinnCompiledObject& compiled_obj,
const CinnLaunchContext& context);

// Set cinn FLAGS (such as FLAGS_cinn_cudnn_deterministic) with paddle's FLAGS.
void SetCinnRuntimeFlags();
} // namespace details

template <typename DeviceContext, typename T>
Expand Down Expand Up @@ -202,7 +206,10 @@ class CinnLaunchOpKernel : public framework::OpKernel<T> {
launch_context->AssignInternalVariable(var_name, tensor);
}

// Step 4. Launch CINN to execute the compiled executable program
// Step 4. Set CINN runtime FLAGS, such as FLAGS_cinn_cudnn_deterministic.
details::SetCinnRuntimeFlags();

// Step 5. Launch CINN to execute the compiled executable program
details::LaunchCinnExecution(cinn_compiled_object, *launch_context);
VLOG(4) << "CinnLaunchOp launch execution done.";
}
Expand Down
5 changes: 3 additions & 2 deletions paddle/fluid/operators/conv_cudnn_op.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -298,11 +298,12 @@ class CUDNNConvOpKernel : public framework::OpKernel<T> {
miopenConvFwdAlgorithm_t algo{};
using search = SearchAlgorithm<miopenConvFwdAlgorithm_t>;
workspace_size = search::GetWorkspaceSize(args);
algo = search::Find<T>(args, exhaustive_search, false, workspace_size, ctx);
algo = search::Find<T>(args, exhaustive_search, deterministic,
workspace_size, ctx);
#else
cudnnConvolutionFwdAlgo_t algo{};
using search = SearchAlgorithm<cudnnConvolutionFwdAlgoPerf_t>;
algo = search::Find<T>(args, exhaustive_search, false, ctx);
algo = search::Find<T>(args, exhaustive_search, deterministic, ctx);
workspace_size = search::GetWorkspaceSize(args, algo);
#endif

Expand Down