[cuDNN][cuDNN V8 API] Always build assuming cuDNN >= 8.0 (#91527)

We've been building with V8 (incl. V8 API) by default for a while now; this PR cleans up some guards for cuDNN < 8.0.

CC @ptrblck @ngimel
Pull Request resolved: https://github.com/pytorch/pytorch/pull/91527
Approved by: https://github.com/ngimel
This commit is contained in:
Eddie Yan
2023-01-13 18:55:37 +00:00
committed by PyTorch MergeBot
parent 4d26903739
commit 4d07ad74f1
23 changed files with 30 additions and 154 deletions

View File

@ -407,6 +407,7 @@ cc_library(
"@cuda//:cusolver",
"@cuda//:nvrtc",
"@cudnn",
"@cudnn_frontend",
],
alwayslink = True,
)

View File

@ -195,9 +195,6 @@ cmake_dependent_option(
cmake_dependent_option(
BUILD_NVFUSER_BENCHMARK "Build C++ binaries for nvfuser benchmarks" OFF
"USE_CUDA" OFF)
cmake_dependent_option(
USE_EXPERIMENTAL_CUDNN_V8_API "Use experimental cuDNN v8 API" ON
"USE_CUDNN" OFF)
option(USE_FBGEMM "Use FBGEMM (quantized 8-bit server operators)" ON)
option(USE_KINETO "Use Kineto profiling library" ON)
option(USE_CUPTI_SO "Use CUPTI as a shared library" ON)

View File

@ -203,6 +203,12 @@ new_local_repository(
path = "/usr/",
)
new_local_repository(
name = "cudnn_frontend",
build_file = "@//third_party:cudnn_frontend.BUILD",
path = "third_party/cudnn_frontend/",
)
local_repository(
name = "com_github_google_flatbuffers",
path = "third_party/flatbuffers",

View File

@ -305,7 +305,6 @@ struct TORCH_CUDA_CPP_API CTCLossDescriptor
void set(cudnnDataType_t datatype) {
AT_CUDNN_CHECK(cudnnSetCTCLossDescriptor(mut_desc(), datatype));
}
#if CUDNN_VERSION >= 7600
void setEx(
cudnnDataType_t datatype,
cudnnLossNormalizationMode_t normMode,
@ -313,7 +312,6 @@ struct TORCH_CUDA_CPP_API CTCLossDescriptor
AT_CUDNN_CHECK(
cudnnSetCTCLossDescriptorEx(mut_desc(), datatype, normMode, gradMode));
}
#endif
};
struct TORCH_CUDA_CPP_API ActivationDescriptor

View File

@ -59,11 +59,7 @@ cudnnBatchNormMode_t getCudnnBatchNormMode(bool training, at::MemoryFormat memor
return CUDNN_BATCHNORM_PER_ACTIVATION;
} else if (training && memory_format == at::MemoryFormat::ChannelsLast) {
#if CUDNN_VERSION >= 7400
return CUDNN_BATCHNORM_SPATIAL_PERSISTENT;
#else
return CUDNN_BATCHNORM_SPATIAL;
#endif // CUDNN_VERSION >= 7400
} else if (training && memory_format == at::MemoryFormat::ChannelsLast3d) {
@ -152,7 +148,6 @@ std::tuple<Tensor, Tensor, Tensor, Tensor> cudnn_batch_norm(
save_mean = at::empty({ num_features }, weight_t.options());
save_var = at::empty({ num_features }, weight_t.options());
#if CUDNN_VERSION >= 7400
auto op = CUDNN_BATCHNORM_OPS_BN;
size_t workspace_size;
AT_CUDNN_CHECK(cudnnGetBatchNormalizationForwardTrainingExWorkspaceSize(
@ -204,22 +199,6 @@ std::tuple<Tensor, Tensor, Tensor, Tensor> cudnn_batch_norm(
workspace_size,
reserve.data_ptr(),
reserve_size));
#else
reserve = at::empty({0}, input->options().dtype(kByte));
AT_CUDNN_CHECK(cudnnBatchNormalizationForwardTraining(
handle, mode, &one, &zero,
idesc.desc(), input->data_ptr(),
idesc.desc(), output->data_ptr(),
wdesc.desc(),
weight->data_ptr(),
bias->data_ptr(),
exponential_average_factor,
at::maybe_data_ptr(running_mean),
at::maybe_data_ptr(running_var),
epsilon,
save_mean.data_ptr(),
save_var.data_ptr()));
#endif // CUDNN_VERSION >= 7400
} else {
reserve = at::empty({0}, input->options().dtype(kByte));
// This keeps a consistent output with native_batch_norm
@ -317,7 +296,6 @@ std::tuple<Tensor, Tensor, Tensor> cudnn_batch_norm_backward(
Constant one(dataType, 1);
Constant zero(dataType, 0);
#if CUDNN_VERSION >= 7400
auto op = CUDNN_BATCHNORM_OPS_BN;
size_t workspace_size;
@ -354,19 +332,6 @@ std::tuple<Tensor, Tensor, Tensor> cudnn_batch_norm_backward(
workspace_size,
reserve->data_ptr(),
reserve->numel()));
#else
AT_CUDNN_CHECK(cudnnBatchNormalizationBackward(
handle, mode, &one, &zero, &one, &zero,
idesc.desc(), input->data_ptr(),
odesc.desc(), grad_output->data_ptr(),
idesc.desc(), grad_input_t.data_ptr(),
wdesc.desc(), weight->data_ptr(),
grad_weight_t.data_ptr(),
grad_bias_t.data_ptr(),
epsilon,
save_mean->data_ptr(),
save_var->data_ptr()));
#endif // CUDNN_VERSION >= 7400
return std::tuple<Tensor,Tensor,Tensor>{grad_input_t, grad_weight_t, grad_bias_t};
}

View File

@ -111,7 +111,6 @@ void raw_cudnn_convolution_add_relu_fallback_out(
#if AT_CUDNN_ENABLED()
#include <ATen/native/cudnn/Macros.h>
#if HAS_CUDNN_V8()
// v7 functions are preserved here to allow for runtime switching to v7
// (e.g., TORCH_CUDNN_V8_API_DISABLED=1).
// Note that v7 forward/backward out can have different behavior from the v8
@ -149,5 +148,4 @@ void raw_cudnn_convolution_add_relu_out_v7(
bool deterministic,
bool allow_tf32);
#endif
#endif
}}

View File

@ -3,7 +3,6 @@
#if AT_CUDNN_ENABLED()
#include <ATen/native/cudnn/Macros.h>
#include <ATen/core/Tensor.h>
#ifndef AT_PER_OPERATOR_HEADERS
@ -60,10 +59,6 @@
// with the best algo, under the hood, cudnn will run with the slower kernel
// since it sees fastest algorithm combination with a sub optimal mathType.
// Note [blocklist fft algorithms for strided dgrad]
// This is a workaround for a CuDNN bug that gave wrong results in certain strided convolution
// gradient setups. Check Issue #16610 for bug details. Bug is there for CUDNN version < 7.5 .
constexpr size_t operator "" _TiB(unsigned long long n) {
return size_t(n) * 1024 * 1024 * 1024 * 1024;
}
@ -225,15 +220,6 @@ size_t getMaxWorkspaceSize(
template<typename perf_t>
std::vector<perf_t> getValidAlgorithms(perf_t *perfResults, const ConvolutionArgs& args, int n_algo) {
// See Note [blocklist fft algorithms for strided dgrad]
#if CUDNN_VERSION < 7500
bool blocklist = std::is_same<decltype(perfResults[0].algo), cudnnConvolutionBwdDataAlgo_t>::value;
int stride_dim = args.input.dim() - 2;
blocklist &= std::any_of(std::begin(args.params.stride),
std::begin(args.params.stride) + stride_dim,
[=](int n){return n != 1;});
#endif
std::vector<perf_t> result;
result.reserve(n_algo);
for (const auto i : c10::irange(n_algo)) {
@ -244,16 +230,6 @@ std::vector<perf_t> getValidAlgorithms(perf_t *perfResults, const ConvolutionArg
if (perf.status == CUDNN_STATUS_SUCCESS) {
if (!args.params.deterministic || perf.determinism == CUDNN_DETERMINISTIC) {
// See Note [blocklist fft algorithms for strided dgrad]
#if CUDNN_VERSION < 7500
bool skip = blocklist;
skip &= (static_cast<cudnnConvolutionBwdDataAlgo_t>(perfResults[i].algo) == CUDNN_CONVOLUTION_BWD_DATA_ALGO_FFT_TILING ||
static_cast<cudnnConvolutionBwdDataAlgo_t>(perfResults[i].algo) == CUDNN_CONVOLUTION_BWD_DATA_ALGO_FFT);
if (skip) {
continue;
}
#endif
result.push_back(perf);
}
}
@ -493,11 +469,9 @@ public:
perfResults[0].mathType = CUDNN_TENSOR_OP_MATH;
} else {
perfResults[0].mathType = CUDNN_DEFAULT_MATH;
#if defined(CUDNN_VERSION) && CUDNN_VERSION >= 8000
if (args.params.dataType == CUDNN_DATA_FLOAT && !args.params.allow_tf32) {
perfResults[0].mathType = CUDNN_FMA_MATH;
}
#endif
}
search::getWorkspaceSize(args, perfResults[0].algo, &(perfResults[0].memory));
return perfResults;
@ -610,14 +584,10 @@ static inline void split_batch_dim_to_32bit_out(
}
#if defined(CUDNN_VERSION) && CUDNN_VERSION >= 8000
#define ASSERT_CORRECT_PRECISION(math_type) \
if (args.params.dataType == CUDNN_DATA_FLOAT) { \
TORCH_INTERNAL_ASSERT(args.params.allow_tf32 || math_type == CUDNN_FMA_MATH); \
}
#else
#define ASSERT_CORRECT_PRECISION(math_type)
#endif // CUDNN_VERSION >= 8000
// ---------------------------------------------------------------------
@ -672,11 +642,7 @@ void raw_cudnn_convolution_forward_out_32bit(
}
#if !HAS_CUDNN_V8()
void raw_cudnn_convolution_forward_out(
#else
void raw_cudnn_convolution_forward_out_v7(
#endif
const Tensor& output, const Tensor& input, const Tensor& weight,
IntArrayRef padding, IntArrayRef stride, IntArrayRef dilation, int64_t groups,
bool benchmark, bool deterministic, bool allow_tf32) {
@ -734,11 +700,7 @@ void raw_cudnn_convolution_backward_input_out_32bit(
);
}
#if !HAS_CUDNN_V8()
void raw_cudnn_convolution_backward_input_out(
#else
void raw_cudnn_convolution_backward_input_out_v7(
#endif
const at::Tensor& grad_input,
const at::Tensor& grad_output,
const at::Tensor& weight,
@ -797,11 +759,7 @@ void raw_cudnn_convolution_backward_weight_out_32bit(
);
}
#if !HAS_CUDNN_V8()
void raw_cudnn_convolution_backward_weight_out(
#else
void raw_cudnn_convolution_backward_weight_out_v7(
#endif
const Tensor& grad_weight, const Tensor& grad_output, const Tensor& input,
IntArrayRef padding, IntArrayRef stride, IntArrayRef dilation, int64_t groups,
bool benchmark, bool deterministic, bool allow_tf32) {
@ -853,12 +811,7 @@ void raw_cudnn_convolution_backward_weight_out_v7(
TORCH_INTERNAL_ASSERT(false, "This case should not be dispatched to cuDNN.");
}
#if !HAS_CUDNN_V8()
void raw_cudnn_convolution_add_relu_out(
#else
void raw_cudnn_convolution_add_relu_out_v7(
#endif
const Tensor& output,
const Tensor& input,
const Tensor& weight,

View File

@ -4,10 +4,6 @@
#if AT_CUDNN_ENABLED()
#include <ATen/native/cudnn/Macros.h>
#if HAS_CUDNN_V8()
#include <ATen/cudnn/cudnn-wrapper.h>
#include <c10/macros/Macros.h>
@ -787,5 +783,4 @@ void raw_cudnn_convolution_add_relu_out(
}} // at::native
#endif // HAS_CUDNN_V8
#endif // AT_CUDNN_ENABLED

View File

@ -18,7 +18,7 @@
#include <ATen/ops/empty_like.h>
#endif
#if (!AT_CUDNN_ENABLED()) || (CUDNN_VERSION < 7600)
#if (!AT_CUDNN_ENABLED())
namespace at { namespace native {

View File

@ -2,8 +2,6 @@
#include <ATen/cuda/CUDAConfig.h> // for the definition of AT_CUDNN_ENABLED
#if AT_CUDNN_ENABLED()
#include <ATen/native/cudnn/Macros.h>
#if HAS_CUDNN_V8()
#include <ATen/core/TensorBase.h>
#include <ATen/core/TensorBody.h>
@ -259,6 +257,5 @@ TORCH_LIBRARY_IMPL(quantized, QuantizedCUDA, m) {
} // namespace native
} // namespace at
#endif // HAS_CUDNN_V8
#endif // AT_CUDNN_ENABLED
#endif // USE_CUDA

View File

@ -3,11 +3,8 @@
#if AT_CUDNN_ENABLED()
#include <ATen/native/cudnn/Macros.h>
#include <c10/util/ArrayRef.h>
#if HAS_CUDNN_V8()
#include <ATen/ATen.h>
#include <ATen/cuda/Exceptions.h>
#include <ATen/cudnn/Handle.h>
@ -432,6 +429,5 @@ TORCH_LIBRARY_IMPL(quantized, QuantizedCUDA, m) {
} // namespace at
#endif // HAS_CUDNN_V8
#endif // AT_CUDNN_ENABLED
#endif // USE_CUDA

View File

@ -3,10 +3,6 @@
#if AT_CUDNN_ENABLED()
#include <ATen/native/cudnn/Macros.h>
#if HAS_CUDNN_V8()
#include <ATen/ATen.h>
#include <torch/library.h>
#include <ATen/native/quantized/cpu/QuantUtils.h>
@ -212,6 +208,5 @@ TORCH_LIBRARY_IMPL(quantized, QuantizedCUDA, m) {
} // namespace native
} // namespace at
#endif // HAS_CUDNN_V8
#endif // AT_CUDNN_ENABLED
#endif // USE_CUDA

View File

@ -3,10 +3,6 @@
#if AT_CUDNN_ENABLED()
#include <ATen/native/cudnn/Macros.h>
#if HAS_CUDNN_V8()
#include <ATen/ATen.h>
#include <ATen/native/quantized/cudnn/utils.h>
#include <ATen/native/quantized/PackedParams.h>
@ -23,6 +19,5 @@ std::tuple<at::Tensor, c10::optional<at::Tensor>> PackedConvWeightCudnn<
template std::tuple<at::Tensor, c10::optional<at::Tensor>> PackedConvWeightCudnn<
2>::unpack();
#endif // HAS_CUDNN_V8
#endif // AT_CUDNN_ENABLED
#endif // USE_CUDA

View File

@ -3,11 +3,8 @@
#if AT_CUDNN_ENABLED()
#include <ATen/native/cudnn/Macros.h>
#include <c10/util/ArrayRef.h>
#if HAS_CUDNN_V8()
#include <ATen/ATen.h>
#include <ATen/cuda/Exceptions.h>
#include <ATen/cudnn/Handle.h>
@ -367,6 +364,5 @@ TORCH_LIBRARY_IMPL(quantized, QuantizedCUDA, m) {
} // namespace at
#endif // HAS_CUDNN_V8
#endif // AT_CUDNN_ENABLED
#endif // USE_CUDA

View File

@ -3,10 +3,6 @@
#if AT_CUDNN_ENABLED()
#include <ATen/native/cudnn/Macros.h>
#if HAS_CUDNN_V8()
#include <ATen/ATen.h>
#include <torch/library.h>
#include <ATen/native/quantized/cudnn/utils.h>
@ -58,6 +54,5 @@ TORCH_LIBRARY_IMPL(quantized, QuantizedCUDA, m) {
} // namespace native
} // namespace at
#endif // HAS_CUDNN_V8
#endif // AT_CUDNN_ENABLED
#endif // USE_CUDA

View File

@ -3,10 +3,6 @@
#if AT_CUDNN_ENABLED()
#include <ATen/native/cudnn/Macros.h>
#if HAS_CUDNN_V8()
#include <ATen/ATen.h>
#include <ATen/native/quantized/cudnn/utils.h>
#include <ATen/native/quantized/PackedParams.h>
@ -18,6 +14,5 @@ std::tuple<at::Tensor, c10::optional<at::Tensor>> PackedLinearWeightCudnn::unpac
return std::tuple<at::Tensor, c10::optional<at::Tensor>>{orig_weight, bias_};
}
#endif // HAS_CUDNN_V8
#endif // AT_CUDNN_ENABLED
#endif // USE_CUDA

View File

@ -3,7 +3,6 @@
#include <ATen/cuda/CUDAConfig.h> // for the definition of AT_CUDNN_ENABLED
#if AT_CUDNN_ENABLED()
#include <ATen/native/cudnn/Macros.h>
#include <ATen/cuda/Exceptions.h>
#include <ATen/cudnn/Descriptors.h>
#include <ATen/cudnn/Handle.h>
@ -54,7 +53,6 @@ Tensor adaptive_avg_pool2d_quantized_cuda(
// TODO: renable these cudnn preprocessors like quantized_max_pool2d_cudnn below when we implement this function with cudnn
#ifdef USE_CUDA
// #if AT_CUDNN_ENABLED()
// #if HAS_CUDNN_V8()
// TODO: limit this to per tensor quantized tensors for now, though should be easy to adapt
// to per channel quantized tensors
TORCH_CHECK(input.qscheme() == at::kPerTensorAffine, "adaptive_avg_pool2d_quantized_cuda oonly supports per tensor quantized tensors");
@ -91,7 +89,6 @@ Tensor quantized_max_pool2d_cudnn(
bool ceil_mode) {
#ifdef USE_CUDA
#if AT_CUDNN_ENABLED()
#if HAS_CUDNN_V8()
check_maxpool2d_params(
kernel_size,
stride,
@ -207,10 +204,6 @@ Tensor quantized_max_pool2d_cudnn(
// recall we casted our input and output to 4D if qx was 3D, so we recast it back to 3D prior to returning
return (ndim == 3 ? qy.view(std::vector<int64_t>(output_shape.begin() + 1, output_shape.end())) : qy);
#else // HAS_CUDNN_V8()
AT_ERROR("at::native::quantized_max_pool2d_cudnn: ATen not compiled with cuDNN v8 support");
return Tensor{}; // never reached, placates the compiler
#endif // HAS_CUDNN_V8()
#else // AT_CUDNN_ENABLED()
AT_ERROR("at::native::quantized_max_pool2d_cudnn: ATen not compiled with cuDNN support");
return Tensor{}; // never reached, placates the compiler

View File

@ -8,10 +8,6 @@ This file contains some of the auxiliary functions used by both Conv.cpp & Linea
#if AT_CUDNN_ENABLED()
#include <ATen/native/cudnn/Macros.h>
#if HAS_CUDNN_V8()
#include <ATen/cudnn/Types.h>
#include <ATen/Tensor.h>
#include <ATen/native/quantized/PackedParams.h>
@ -354,6 +350,5 @@ cudnn_frontend::ExecutionPlan get_execplan_from_heuristics_else_fall_back(cudnn_
} // anonymous
} // cudnn_utils
#endif // HAS_CUDNN_V8
#endif // AT_CUDNN_ENABLED
#endif // USE_CUDA

View File

@ -1393,12 +1393,6 @@ elseif(USE_ROCM)
target_compile_definitions(torch_hip PRIVATE "-DTORCH_HIP_BUILD_MAIN_LIB")
endif()
if(USE_EXPERIMENTAL_CUDNN_V8_API)
if(USE_CUDA)
target_compile_definitions(torch_cuda PRIVATE "-DUSE_EXPERIMENTAL_CUDNN_V8_API")
endif()
endif()
set(EXPERIMENTAL_SINGLE_THREAD_POOL "0" CACHE STRING
"Experimental option to use a single thread pool for inter- and intra-op parallelism")
if("${EXPERIMENTAL_SINGLE_THREAD_POOL}")

View File

@ -77,7 +77,6 @@ function(caffe2_print_configuration_summary)
message(STATUS " Split CUDA : ${BUILD_SPLIT_CUDA}")
message(STATUS " CUDA static link : ${CAFFE2_STATIC_LINK_CUDA}")
message(STATUS " USE_CUDNN : ${USE_CUDNN}")
message(STATUS " USE_EXPERIMENTAL_CUDNN_V8_API: ${USE_EXPERIMENTAL_CUDNN_V8_API}")
message(STATUS " CUDA version : ${CUDA_VERSION}")
message(STATUS " USE_FLASH_ATTENTION : ${USE_FLASH_ATTENTION}")
if(${USE_CUDNN})

22
third_party/cudnn_frontend.BUILD vendored Normal file
View File

@ -0,0 +1,22 @@
# Adopted from: https://github.com/tensorflow/tensorflow/blob/master/third_party/cudnn_frontend.BUILD
# Description:
# The cuDNN Frontend API is a C++ header-only library that demonstrates how
# to use the cuDNN C backend API.
load("@rules_cc//cc:defs.bzl", "cc_library")
package(
default_visibility = ["//visibility:public"],
)
licenses(["notice"]) # MIT
exports_files(["LICENSE.txt"])
cc_library(
name = "cudnn_frontend",
hdrs = glob(["include/**"]),
includes = ["include/"],
include_prefix = "third_party/cudnn_frontend",
)

View File

@ -150,10 +150,6 @@ if(USE_ROCM)
list(APPEND TORCH_PYTHON_INCLUDE_DIRECTORIES ${roctracer_INCLUDE_DIRS})
endif()
if(USE_EXPERIMENTAL_CUDNN_V8_API)
list(APPEND TORCH_PYTHON_COMPILE_DEFINITIONS USE_EXPERIMENTAL_CUDNN_V8_API)
endif()
if(USE_CUDNN OR USE_ROCM)
list(APPEND TORCH_PYTHON_SRCS
${TORCH_SRC_DIR}/csrc/cuda/shared/cudnn.cpp

View File

@ -1146,12 +1146,7 @@ PyObject* THCPModule_setBenchmarkLimitCuDNN(PyObject* _unused, PyObject* arg) {
"cuDNN Benchmark limit is not supported in MIOpen and will have no effect.");
#endif
#if AT_CUDNN_ENABLED()
#if HAS_CUDNN_V8()
at::globalContext().setBenchmarkLimitCuDNN(benchmark_limit);
#else
TORCH_WARN_ONCE(
"cuDNN Benchmark limit is not supported with cuDNN v7 API and will have no effect.");
#endif
#endif
Py_RETURN_NONE;
}