[arch-commits] Commit in root/trunk (5 files)

Konstantin Gizdov kgizdov at archlinux.org
Thu Jul 23 20:24:10 UTC 2020


    Date: Thursday, July 23, 2020 @ 20:24:09
  Author: kgizdov
Revision: 665090

upgpkg: root 6.22.00-1

Added:
  root/trunk/adapt_tmva_to_support_cudnn8.patch
Modified:
  root/trunk/PKGBUILD
  root/trunk/nbman-for-arch.patch
Deleted:
  root/trunk/add_missing_include_ROOT-10315.patch
  root/trunk/fix_relative_paths.patch

--------------------------------------+
 PKGBUILD                             |   20 
 adapt_tmva_to_support_cudnn8.patch   | 1130 +++++++++++++++++++++++++++++++++
 add_missing_include_ROOT-10315.patch |   22 
 fix_relative_paths.patch             |   39 -
 nbman-for-arch.patch                 |   17 
 5 files changed, 1148 insertions(+), 80 deletions(-)

Modified: PKGBUILD
===================================================================
--- PKGBUILD	2020-07-23 20:13:52 UTC (rev 665089)
+++ PKGBUILD	2020-07-23 20:24:09 UTC (rev 665090)
@@ -6,7 +6,7 @@
 
 pkgbase=root
 pkgname=('root' 'root-cuda')
-pkgver=6.20.06
+pkgver=6.22.00
 pkgrel=1
 pkgdesc='C++ data analysis framework and interpreter from CERN'
 arch=('x86_64')
@@ -21,7 +21,7 @@
     'cuda'
     'cudnn'
     'gcc-fortran'
-    'gcc8-fortran'
+    'gcc9-fortran'
     'git'
     'go'
     'libxml2'
@@ -90,12 +90,11 @@
     'settings.cmake'
     'settings-cuda.cmake'
     'jupyter_notebook_config.py'
-    'add_missing_include_ROOT-10315.patch'
-    'fix_relative_paths.patch'
     'nbman-for-arch.patch'
     'thisroot.fail'
+    'adapt_tmva_to_support_cudnn8.patch'
 )
-sha512sums=('232fd5253e83eb02ad33d03941c7c83d39d5f6b0162bd42594e44a0c1c08deade42ae5793e571db767ce4fa0a582aa332b9d1869e621b825d1eb1162819c98c6'
+sha512sums=('9e3c54bbc146b0abb0a2d960af380255ec59d0b3a11a4a97a2a25cb7ac567b07280c4eb48dddf99c1fa2e692881f6396a842ce125d3a253037e52f719739f01e'
             'af8f178fc9df66997d5495b271e38adcd1636aab4c8fc994c6600c2496127829d831250d73d3fc229b02dfe49b9867d0be979beacb959f2f3a05351b8118a4a6'
             '1fe6f4aa09d583d33f27cc766f4935510bb7ab6bbb8d4700baa1aaab92ea6c876500b67da1e4f6e0b510aa5616e4e193b860264b86925de85f2d9f558d75d5dc'
             '3c81d255a17b902ffac0187af1752847036137e16641a88b17eef0d9c944e6f0d3c954bc93307d6270603f43f6c23f2e04f98dc7a68f9d076dbaa8006a2527d6'
@@ -102,10 +101,9 @@
             '9ee5b6606dbd352608a2a4998344ca4026d677c86823e62fff615f6e84efcecdffc07a1e9182a356aa35035e7f35df5a107127722a6bad4b97d1f49cffebf5b9'
             '7665bc8cbe79162e0b969b08802e1b7b2ed22ed8b1402d50cf194172a644f647dcaf0f5abb76f8b6007dfab8dbc811604479be826b345d8fd77edfb51032110b'
             '1c905ee7a3f8f5f3f567d957f9be6b503a8631565d4d9b9bfea5e496ef86865c5a8be1a1f8c7842754029879cf0afd2465249f532a116cc43660aa2e460ae682'
-            '19b46d3b6a8083d9461f5d4f57845a1145b80de699e8b773cee7edecad0b064fe570faffa226720078273f40a88fe8ba933b7d944925391ad869f0b9cdbf9579'
-            'fa2f3b9b8717bfd11b6743a4457bbff104e547112bcec864f212510b1954a92809d9bde5ada8702d11659a150f2957daed4e4d2621858b0bc4e234d9f6924244'
-            'f4152ecddbef8079e7507625c3cfb2dcd6d3c844f38827f293d7c1915ff9be5ea3e4414080b51a83f86128b954520d8ec113349a5a07ba39fc87047b72bf0136'
-            'ff555ac4db568affe139701907f86d919a2206f3e304f69dd317b756ea0904b5934d9364a524060778aa507809ce78448621619bb34039ba34c5a71af71a4a8c')
+            '12814f50b7016bd86d3f91e0e31c052783a0c0fa72b7d6a072d3ae6f86c2437323d585e531235377ebbfdd9cb76abd7da84d9631de821151547f1d4b13417e69'
+            'ff555ac4db568affe139701907f86d919a2206f3e304f69dd317b756ea0904b5934d9364a524060778aa507809ce78448621619bb34039ba34c5a71af71a4a8c'
+            '2ae126795df4127c27a6287a1499bdb8b2bacb74cfbec17dabe378a5fb9fc7c755644e4090a4da1d0045bf5d4f542f06da827a0f48a5927ee8509874045f18b6')
 
 get_pyver () {
     python -c 'import sys; print(str(sys.version_info[0]) + "." + str(sys.version_info[1]))'
@@ -167,8 +165,8 @@
     mkdir -p "${srcdir}/build-cuda"
     cd "${srcdir}/build-cuda"
 
-    CC=/usr/bin/gcc-8 \
-    CXX=/usr/bin/g++-8 \
+    CC=/usr/bin/gcc-9 \
+    CXX=/usr/bin/g++-9 \
     cmake -C "${srcdir}/settings-cuda.cmake" \
         ${CUSTOM_CMAKE_FLAGS} \
         "${srcdir}/${pkgbase}-${pkgver}-cuda"

Added: adapt_tmva_to_support_cudnn8.patch
===================================================================
--- adapt_tmva_to_support_cudnn8.patch	                        (rev 0)
+++ adapt_tmva_to_support_cudnn8.patch	2020-07-23 20:24:09 UTC (rev 665090)
@@ -0,0 +1,1130 @@
+From 05739e6b01fb34b5ef40e1a584107876e68e4b77 Mon Sep 17 00:00:00 2001
+From: Konstantin Gizdov <kgizdov at gmail.com>
+Date: Tue, 21 Jul 2020 15:13:57 +0300
+Subject: [PATCH 01/10] update deprecated function call name to backward
+ compatible one
+
+---
+ tmva/tmva/src/DNN/Architectures/Cudnn/RecurrentPropagation.cu | 4 ++++
+ 1 file changed, 4 insertions(+)
+
+diff --git a/tmva/tmva/src/DNN/Architectures/Cudnn/RecurrentPropagation.cu b/tmva/tmva/src/DNN/Architectures/Cudnn/RecurrentPropagation.cu
+index 058cee28424..60289ec2fdd 100644
+--- a/tmva/tmva/src/DNN/Architectures/Cudnn/RecurrentPropagation.cu
++++ b/tmva/tmva/src/DNN/Architectures/Cudnn/RecurrentPropagation.cu
+@@ -132,7 +132,11 @@ void TCudnn<AFloat>::InitializeRecurrentDescriptors(TDescriptors *&descriptors,
+    cudnnDataType_t mathPrec = CUDNN_DATA_FLOAT;
+    if      (std::is_same<AFloat, double>::value) { mathPrec = CUDNN_DATA_DOUBLE;}
+ 
++#if (CUDNN_VERSION >= 8000)
++   CUDNNCHECK(cudnnSetRNNDescriptor_v6(handle, rnnDescriptors->LayerDescriptor, hiddenSize, numLayers, rnnDescriptors->HelperDescriptor,
++#else
+    CUDNNCHECK(cudnnSetRNNDescriptor(handle, rnnDescriptors->LayerDescriptor, hiddenSize, numLayers, rnnDescriptors->HelperDescriptor,
++#endif
+       inputMode, direction, mode, algo, mathPrec) );
+ 
+ 
+
+From 90baa4f6ad10076fa148f5aa06ef432bd0f34208 Mon Sep 17 00:00:00 2001
+From: Konstantin Gizdov <kgizdov at gmail.com>
+Date: Tue, 21 Jul 2020 19:06:09 +0300
+Subject: [PATCH 02/10] adapt convolution forward to cuDNN 8
+
+---
+ .../src/DNN/Architectures/Cudnn/Propagate.cu  | 77 ++++++++++++++++++-
+ 1 file changed, 76 insertions(+), 1 deletion(-)
+
+diff --git a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+index 7a57b6bf104..cc953ee45f9 100644
+--- a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
++++ b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+@@ -27,6 +27,9 @@
+ // #include "Kernels.cuh"*/
+ // #include <math.h>
+ 
++// for std::numeric_limits<T>::max()
++#include <limits>
++
+ namespace TMVA {
+ namespace DNN  {
+ 
+@@ -378,7 +381,78 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    cudnnHandle_t cudnnHandle = outputTensor.GetCudnnHandle();
+ 
+    // cuDNN decides which algorithm to use
+-   // More detailed alternative: cudnnFindConvolutionForwardAlgorithm
++#if (CUDNN_VERSION >= 8000)
++   /**
++    * I'm sure there may be a faster way, but this works
++    */
++   int convRequestedAlgoCount{8}; // requestedAlgoCount is setting how many algorithms to try, can be tuned, fixed for now as all available
++   cudnnConvolutionDescriptor_t tempConvDescriptor;
++   CUDDNCHECK(cudnnCreateConvolutionDescriptor(&tempConvDescriptor));
++   cudnnTensorDescriptor_t  outputTensorDescriptor;
++   CUDNNCHECK(cudnnCreateTensorDescriptor(&outputTensorDescriptor));
++   CUDNNCHECK(cudnnSetTensor4dDescriptor(outputTensorDescriptor,
++                                             CUDNN_TENSOR_NCHW,  // Layout of the tensor in memory
++                                             Tensor_t::GetDataType(),
++                                             (int)L->GetBatchSize(),
++                                             (int)L->GetDepth(),
++                                             (int)L->GetHeight(),
++                                             (int)L->GetWidth()));
++   int algoCount;
++   cudnnConvolutionFwdAlgoPerf_t convPerfResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
++   CUDNNCHECK(cudnnFindConvolutionForwardAlgorithm(
++      cudnnHandle,
++      inputTensorDescriptor,
++      convDescriptors->WeightsDescriptor,
++      tempConvDescriptor,
++      outputTensorDescriptor,
++      convRequestedAlgoCount,
++      &algoCount,
++      &convPerfResults));
++   // we could also do it with the expert mode (cudnnFindConvolutionForwardAlgorithmEx),
++   // but we arrive at an chicken or egg problem:
++   // workspace size is calculated from chosen forward algorithm,
++   // but finding a forward algorithm depends on workspace size...
++   // i.e.
++   // Tensor_t & inputTensor = L->GetInput();
++   // inputTensor = Tensor_t(inputTensor.GetDeviceBuffer(),{ L->GetBatchSize(), L->GetInputDepth(), L->GetInputHeight(), L->GetInputWidth() },GetTensorLayout(),0,0);
++   // CUDNNCHECK(cudnnFindConvolutionForwardAlgorithmEx(
++   //    cudnnHandle,
++   //    inputTensorDescriptor,
++   //    &inputTensor,
++   //    convDescriptors->WeightsDescriptor,
++   //    &filters,
++   //    tempConvDescriptor,
++   //    outputTensorDescriptor,
++   //    &outputTensor,
++   //    convRequestedAlgoCount,
++   //    &algoCount,
++   //    &convPerfResults,
++   //    &convWorkspace,
++   //    convWorkspace->ForwardWorkspaceSize));
++   // instead choose either fastest or lowest memory algo as per preference
++   int algoIdx{0};
++   if (CNNOptions::ConvMaxWorkspaceSize != 0) {  // prefer fastest
++      float temp_runtime{std::numeric_limits<float>::max()};
++      for (int i = 0; i < algoCount; ++i) {
++         if (convPerfResults[i].status != 0) continue;
++         if (convPerfResults[i].time < temp_runtime) {
++            temp_runtime = convPerfResults[i].time;
++            algoIdx = i;
++         }
++      }
++   } else {  // prefer smallest workspace size
++      size_t temp_memsize{std::numeric_limits<size_t>::max()};
++      for (int i = 0; i < algoCount; ++i) {
++         if (convPerfResults[i].status != 0) continue;
++         if (convPerfResults[i].memory < temp_memsize) {
++            temp_memsize = convPerfResults[i].memory;
++            algoIdx = i;
++         }
++      }
++   }
++   convWorkspace->AlgorithmForward = convPerfResults[algoIdx].algo;
++#else
++   // More detailed alternative: cudnnFindConvolutionForwardAlgorithm (only option in newer cuDNN versions)
+    cudnnConvolutionFwdPreference_t preferenceFwd = (CNNOptions::ConvMaxWorkspaceSize !=0) ? CUDNN_CONVOLUTION_FWD_PREFER_FASTEST :
+                                                    CUDNN_CONVOLUTION_FWD_NO_WORKSPACE;
+ 
+@@ -389,6 +463,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+       outputTensor.GetTensorDescriptor(), preferenceFwd,
+       memLimit, // Memory limit in bytes for mode CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT
+       &convWorkspace->AlgorithmForward));
++#endif
+ 
+    // Allocate memory for the convolution
+    //size_t workSpaceSizeInBytes = 0;
+
+From d9b5e2f82917e7183b9f45a49135641981741477 Mon Sep 17 00:00:00 2001
+From: Konstantin Gizdov <kgizdov at gmail.com>
+Date: Tue, 21 Jul 2020 19:34:00 +0300
+Subject: [PATCH 03/10] adapt convolution backward to cuDNN 8
+
+---
+ .../src/DNN/Architectures/Cudnn/Propagate.cu  | 72 +++++++++++++++++++
+ 1 file changed, 72 insertions(+)
+
+diff --git a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+index cc953ee45f9..85a5c3aa175 100644
+--- a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
++++ b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+@@ -515,6 +515,77 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    // dx : Activation gradient to be computed                               -> activationGradients [in place op]
+    // dy : Gradient of activation from the following layer (backpropagation)-> activationGradients
+ 
++#if (CUDNN_VERSION >= 8000)
++   /**
++    * I'm sure there may be a faster way, but this works
++    */
++   convRequestedAlgoCount = 6; // reset to max number of available backward algorithms
++   cudnnConvolutionDescriptor_t tempConvBwdDescriptor;
++   CUDDNCHECK(cudnnCreateConvolutionDescriptor(&tempConvBwdDescriptor));
++   cudnnTensorDescriptor_t  outputBwdTensorDescriptor;
++   CUDNNCHECK(cudnnCreateTensorDescriptor(&outputBwdTensorDescriptor));
++   CUDNNCHECK(cudnnSetTensor4dDescriptor(outputBwdTensorDescriptor,
++                                             CUDNN_TENSOR_NCHW,  // Layout of the tensor in memory
++                                             Tensor_t::GetDataType(),
++                                             (int)L->GetBatchSize(),
++                                             (int)L->GetInputDepth(),
++                                             (int)L->GetInputHeight(),
++                                             (int)L->GetInputWidth()));
++   int algoCount;
++   cudnnConvolutionBwdDataAlgoPerf_t convPerfBwdResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
++   CUDNNCHECK(cudnnFindConvolutionBackwardDataAlgorithm(
++      cudnnHandle,
++      convDescriptors->WeightsDescriptor,
++      activationGradientsBackwardDescriptor,
++      tempConvBwdDescriptor,
++      outputBwdTensorDescriptor,
++      convRequestedAlgoCount,
++      &algoCount,
++      &convPerfBwdResults));
++   // we could also do it with the expert mode (cudnnFindConvolutionForwardAlgorithmEx),
++   // but we arrive at an chicken or egg problem:
++   // workspace size is calculated from chosen forward algorithm,
++   // but finding a forward algorithm depends on workspace size...
++   // i.e.
++   // Tensor_t & outputBwdTensor = L->GetInput();
++   // outputBwdTensor = Tensor_t(outputBwdTensor.GetDeviceBuffer(),{ L->GetBatchSize(), L->GetInputDepth(), L->GetInputHeight(), L->GetInputWidth() },GetTensorLayout(),0,0);
++   // CUDNNCHECK(cudnnFindConvolutionBackwardDataAlgorithmEx(
++   //    cudnnHandle,
++   //    convDescriptors->WeightsDescriptor,
++   //    &filters,
++   //    activationGradientsBackwardDescriptor,
++   //    &activationGradientsBackwardTensor,
++   //    tempConvBwdDescriptor,
++   //    outputBwdTensorDescriptor,
++   //    &outputBwdTensor,
++   //    convRequestedAlgoCount,
++   //    &algoCount,
++   //    &convPerfBwdResults,
++   //    &convWorkspace,
++   //    convWorkspace->ForwardWorkspaceSize));
++   // instead choose either fastest or lowest memory algo as per preference
++   int algoIdx{0};
++   if (CNNOptions::ConvMaxWorkspaceSize != 0) {  // prefer fastest
++      float temp_runtime{std::numeric_limits<float>::max()};
++      for (int i = 0; i < algoCount; ++i) {
++         if (convPerfBwdResults[i].status != 0) continue;
++         if (convPerfBwdResults[i].time < temp_runtime) {
++            temp_runtime = convPerfBwdResults[i].time;
++            algoIdx = i;
++         }
++      }
++   } else {  // prefer smallest workspace size
++      size_t temp_memsize{std::numeric_limits<size_t>::max()};
++      for (int i = 0; i < algoCount; ++i) {
++         if (convPerfBwdResults[i].status != 0) continue;
++         if (convPerfBwdResults[i].memory < temp_memsize) {
++            temp_memsize = convPerfBwdResults[i].memory;
++            algoIdx = i;
++         }
++      }
++   }
++   convWorkspace->AlgorithmBackward = convPerfBwdResults[algoIdx].algo;
++#else
+    cudnnConvolutionBwdDataPreference_t preferenceBwdData =
+       (CNNOptions::ConvMaxWorkspaceSize != 0) ? CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST : CUDNN_CONVOLUTION_BWD_DATA_NO_WORKSPACE;
+ 
+@@ -525,6 +596,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+                                                       activationGradientsBackwardDescriptor,
+                                                       preferenceBwdData, memLimit,
+                                                       &convWorkspace->AlgorithmBackward));
++#endif
+ 
+    std::cout << "CONV BWD Data Algo used  is "  << convWorkspace->AlgorithmBackward << std::endl;
+    //CUDNNCHECK(cudnnSetConvolutionMathType(convDescriptors->LayerDescriptor, CUDNN_TENSOR_OP_MATH));
+
+From 526b7177c0201be1d0c6b36de0772b7d2ecb90d5 Mon Sep 17 00:00:00 2001
+From: Konstantin Gizdov <kgizdov at gmail.com>
+Date: Wed, 22 Jul 2020 11:50:29 +0300
+Subject: [PATCH 04/10] fix typo and re-declarations
+
+---
+ tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu | 11 +++++------
+ 1 file changed, 5 insertions(+), 6 deletions(-)
+
+diff --git a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+index 85a5c3aa175..1b7e3e845d8 100644
+--- a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
++++ b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+@@ -387,7 +387,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+     */
+    int convRequestedAlgoCount{8}; // requestedAlgoCount is setting how many algorithms to try, can be tuned, fixed for now as all available
+    cudnnConvolutionDescriptor_t tempConvDescriptor;
+-   CUDDNCHECK(cudnnCreateConvolutionDescriptor(&tempConvDescriptor));
++   CUDNNCHECK(cudnnCreateConvolutionDescriptor(&tempConvDescriptor));
+    cudnnTensorDescriptor_t  outputTensorDescriptor;
+    CUDNNCHECK(cudnnCreateTensorDescriptor(&outputTensorDescriptor));
+    CUDNNCHECK(cudnnSetTensor4dDescriptor(outputTensorDescriptor,
+@@ -407,7 +407,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+       outputTensorDescriptor,
+       convRequestedAlgoCount,
+       &algoCount,
+-      &convPerfResults));
++      convPerfResults));
+    // we could also do it with the expert mode (cudnnFindConvolutionForwardAlgorithmEx),
+    // but we arrive at an chicken or egg problem:
+    // workspace size is calculated from chosen forward algorithm,
+@@ -521,7 +521,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+     */
+    convRequestedAlgoCount = 6; // reset to max number of available backward algorithms
+    cudnnConvolutionDescriptor_t tempConvBwdDescriptor;
+-   CUDDNCHECK(cudnnCreateConvolutionDescriptor(&tempConvBwdDescriptor));
++   CUDNNCHECK(cudnnCreateConvolutionDescriptor(&tempConvBwdDescriptor));
+    cudnnTensorDescriptor_t  outputBwdTensorDescriptor;
+    CUDNNCHECK(cudnnCreateTensorDescriptor(&outputBwdTensorDescriptor));
+    CUDNNCHECK(cudnnSetTensor4dDescriptor(outputBwdTensorDescriptor,
+@@ -531,7 +531,6 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+                                              (int)L->GetInputDepth(),
+                                              (int)L->GetInputHeight(),
+                                              (int)L->GetInputWidth()));
+-   int algoCount;
+    cudnnConvolutionBwdDataAlgoPerf_t convPerfBwdResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
+    CUDNNCHECK(cudnnFindConvolutionBackwardDataAlgorithm(
+       cudnnHandle,
+@@ -541,7 +540,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+       outputBwdTensorDescriptor,
+       convRequestedAlgoCount,
+       &algoCount,
+-      &convPerfBwdResults));
++      convPerfBwdResults));
+    // we could also do it with the expert mode (cudnnFindConvolutionForwardAlgorithmEx),
+    // but we arrive at an chicken or egg problem:
+    // workspace size is calculated from chosen forward algorithm,
+@@ -564,7 +563,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    //    &convWorkspace,
+    //    convWorkspace->ForwardWorkspaceSize));
+    // instead choose either fastest or lowest memory algo as per preference
+-   int algoIdx{0};
++   algoIdx = 0;
+    if (CNNOptions::ConvMaxWorkspaceSize != 0) {  // prefer fastest
+       float temp_runtime{std::numeric_limits<float>::max()};
+       for (int i = 0; i < algoCount; ++i) {
+
+From 6d84e765322a72c48de00b4a9b7471da8a15fece Mon Sep 17 00:00:00 2001
+From: Konstantin Gizdov <kgizdov at gmail.com>
+Date: Wed, 22 Jul 2020 17:00:01 +0300
+Subject: [PATCH 05/10] implement workspace limits, fix an algoruthm preference
+ bug and rewrite relevant sections
+
+---
+ .../src/DNN/Architectures/Cudnn/Propagate.cu  | 273 ++++++++++--------
+ 1 file changed, 151 insertions(+), 122 deletions(-)
+
+diff --git a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+index 1b7e3e845d8..2049e2b9195 100644
+--- a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
++++ b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+@@ -333,35 +333,108 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+                                              TDescriptors * & descriptors,
+                                              const DNN::CNN::TConvParams & /*params*/,
+                                              ConvLayer_t *L) {
+-   auto convWorkspace = new ConvWorkspace_t ();
++   auto convWorkspace = new ConvWorkspace_t();
++   size_t memLimit = (CNNOptions::ConvMaxWorkspaceSize > 0) ? static_cast<size_t>(CNNOptions::ConvMaxWorkspaceSize) : 0;
+    auto convDescriptors = static_cast<ConvDescriptors_t *>(descriptors);
++   // can we do the following and substitute below???
++   // auto weightsDescriptor{convDescriptors->WeightsDescriptor};
++   // auto convDescriptor{convDescriptors->LayerDescriptor};
+ 
++#if (CUDNN_VERSION >= 8000)
++   enum algoPreference { no_workspace, fastest, workspace_limit };
++   algoPreference algoChoice;
++   auto choose_algo = [](algoPreference const& algoPref, auto&& perfResults, size_t memLim = std::numeric_limits<size_t>::max()) -> int {
++      int algoIdx{0};
++      if (algoPref == algoPreference::fastest) {  // prefer fastest
++         float temp_runtime{std::numeric_limits<float>::max()};
++         for (int i = 0; i < algoCount; ++i) {
++            if (PerfResults[i].status == CUDNN_STATUS_SUCCESS && PerfResults[i].time < temp_runtime) {
++               temp_runtime = PerfResults[i].time;
++               algoIdx = i;
++            }
++         }
++      } else if (algoPref == algoPreference::workspace_limit) {  // constrain to workspace size
++         float temp_runtime{std::numeric_limits<float>::max()};
++         for (int i = 0; i < algoCount; ++i) {
++            if (PerfResults[i].status == CUDNN_STATUS_SUCCESS && PerfResults[i].time < temp_runtime && PerfResults[i].memory <= memLim) {
++               temp_runtime = PerfResults[i].time;
++               algoIdx = i;
++            }
++         }
++      } else {  // prefer smallest workspace size
++         size_t temp_memsize{std::numeric_limits<size_t>::max()};
++         for (int i = 0; i < algoCount; ++i) {
++            if (PerfResults[i].status == CUDNN_STATUS_SUCCESS && PerfResults[i].memory < temp_memsize) {
++               temp_memsize = PerfResults[i].memory;
++               algoIdx = i;
++            }
++         }
++      }
++      return algoIdx;
++   };
++#else
++   // More detailed alternative: cudnnFindConvolutionForwardAlgorithm (only option in newer cuDNN versions)
++   cudnnConvolutionFwdPreference_t preferenceFwd;
++   cudnnConvolutionBwdDataPreference_t preferenceBwdData;
++   cudnnConvolutionBwdFilterPreference_t preferenceBwdFilter;
++#endif
++   // decide on algorithm preference early
++   if (CNNOptions::ConvMaxWorkspaceSize < 0) {
++      // no workspace case
++#if (CUDNN_VERSION >= 8000)
++      algoChoice = no_workspace;
++#else
++      preferenceFwd = CUDNN_CONVOLUTION_FWD_NO_WORKSPACE;
++      preferenceBwdData = CUDNN_CONVOLUTION_BWD_DATA_NO_WORKSPACE;
++      preferenceBwdFilter = CUDNN_CONVOLUTION_BWD_FILTER_NO_WORKSPACE;
++#endif
++
++   } else if (CNNOptions::ConvMaxWorkspaceSize == 0) {
++      // fastest overall
++#if (CUDNN_VERSION >= 8000)
++      algoChoice = fastest;
++#else
++      preferenceFwd = CUDNN_CONVOLUTION_FWD_PREFER_FASTEST;
++      preferenceBwdData = CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST;
++      preferenceBwdFilter = CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST;
++#endif
++
++   } else {
++      // fastest in memory limit
++#if (CUDNN_VERSION >= 8000)
++      algoChoice = workspace_limit;
++#else
++      preferenceFwd = CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT;
++      preferenceBwdData = CUDNN_CONVOLUTION_BWD_DATA_SPECIFY_WORKSPACE_LIMIT;
++      preferenceBwdFilter = CUDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT;
++#endif
++   }
+    // fix the weight tensor shapes
+    // by default the weights are columnmajor, set them to be row major . At this points
+    // they are not yet initialized
+    Tensor_t & filters = L->GetWeightsAt(0);
+-   filters = Tensor_t (filters.GetDeviceBuffer(), {L->GetDepth(),L->GetInputDepth(), L->GetFilterHeight(),L->GetFilterWidth()}, MemoryLayout::RowMajor, 0, 0 );
+-   //PrintTensor(L->GetWeightsAt(0));
++   filters = Tensor_t(filters.GetDeviceBuffer(), {L->GetDepth(), L->GetInputDepth(), L->GetFilterHeight(), L->GetFilterWidth()}, MemoryLayout::RowMajor, 0, 0);
++   // PrintTensor(L->GetWeightsAt(0));
+    Tensor_t & biases = L->GetBiasesAt(0);
+-   biases = Tensor_t (biases.GetDeviceBuffer(), {1, L->GetDepth(),1,1}, GetTensorLayout(), 0, 0 );
++   biases = Tensor_t(biases.GetDeviceBuffer(), {1, L->GetDepth(), 1, 1}, GetTensorLayout(), 0, 0);
+ 
+    Tensor_t & outputTensor = L->GetOutput();
+-   outputTensor = Tensor_t(outputTensor.GetDeviceBuffer(),{ L->GetBatchSize(), L->GetDepth(), L->GetHeight(), L->GetWidth() },GetTensorLayout(),0,0 );
++   outputTensor = Tensor_t(outputTensor.GetDeviceBuffer(), {L->GetBatchSize(), L->GetDepth(), L->GetHeight(), L->GetWidth()}, GetTensorLayout(), 0, 0);
+    Tensor_t & inputActivation = L->GetInputActivation();
+-   inputActivation = Tensor_t(inputActivation.GetDeviceBuffer(),outputTensor.GetShape() ,GetTensorLayout(),0,0 );
++   inputActivation = Tensor_t(inputActivation.GetDeviceBuffer(),outputTensor.GetShape() ,GetTensorLayout(), 0, 0);
+ 
+    Tensor_t &  activationGradients = L->GetActivationGradients();
+-   activationGradients =  Tensor_t(activationGradients.GetDeviceBuffer(),outputTensor.GetShape() ,GetTensorLayout(),0,0 );
++   activationGradients = Tensor_t(activationGradients.GetDeviceBuffer(),outputTensor.GetShape(), GetTensorLayout(), 0, 0);
+ 
+    Tensor_t & weightGradients = L->GetWeightGradientsAt(0);
+-   weightGradients = Tensor_t( weightGradients.GetDeviceBuffer(), filters.GetShape(), GetTensorLayout(), 0, 0 );
++   weightGradients = Tensor_t(weightGradients.GetDeviceBuffer(), filters.GetShape(), GetTensorLayout(), 0, 0);
+ 
+    Tensor_t & biasGradients = L->GetBiasGradientsAt(0);
+-   biasGradients = Tensor_t( biasGradients.GetDeviceBuffer(), biases.GetShape(), GetTensorLayout(), 0, 0 );
++   biasGradients = Tensor_t(biasGradients.GetDeviceBuffer(), biases.GetShape(), GetTensorLayout(), 0, 0);
+ 
+ 
+    // FIXME: Use descriptors instead (Tensor device memory is otherwise allocated during initialization)
+-   //Tensor_t inputTensor  ({L->GetBatchSize(), L->GetInputDepth(), L->GetInputHeight(), L->GetInputWidth()}, MemoryLayout::RowMajor, 0, 0);
++   // Tensor_t inputTensor  ({L->GetBatchSize(), L->GetInputDepth(), L->GetInputHeight(), L->GetInputWidth()}, MemoryLayout::RowMajor, 0, 0);
+    cudnnTensorDescriptor_t  inputTensorDescriptor;
+    CUDNNCHECK(cudnnCreateTensorDescriptor(&inputTensorDescriptor) );
+    CUDNNCHECK(cudnnSetTensor4dDescriptor(inputTensorDescriptor,
+@@ -385,79 +458,44 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    /**
+     * I'm sure there may be a faster way, but this works
+     */
+-   int convRequestedAlgoCount{8}; // requestedAlgoCount is setting how many algorithms to try, can be tuned, fixed for now as all available
+-   cudnnConvolutionDescriptor_t tempConvDescriptor;
+-   CUDNNCHECK(cudnnCreateConvolutionDescriptor(&tempConvDescriptor));
+-   cudnnTensorDescriptor_t  outputTensorDescriptor;
+-   CUDNNCHECK(cudnnCreateTensorDescriptor(&outputTensorDescriptor));
+-   CUDNNCHECK(cudnnSetTensor4dDescriptor(outputTensorDescriptor,
+-                                             CUDNN_TENSOR_NCHW,  // Layout of the tensor in memory
+-                                             Tensor_t::GetDataType(),
+-                                             (int)L->GetBatchSize(),
+-                                             (int)L->GetDepth(),
+-                                             (int)L->GetHeight(),
+-                                             (int)L->GetWidth()));
++   int convRequestedAlgoCount{8};  // requestedAlgoCount is setting how many algorithms to try, can be tuned, fixed for now as all available
++
+    int algoCount;
+    cudnnConvolutionFwdAlgoPerf_t convPerfResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
+-   CUDNNCHECK(cudnnFindConvolutionForwardAlgorithm(
+-      cudnnHandle,
+-      inputTensorDescriptor,
+-      convDescriptors->WeightsDescriptor,
+-      tempConvDescriptor,
+-      outputTensorDescriptor,
+-      convRequestedAlgoCount,
+-      &algoCount,
+-      convPerfResults));
++   CUDNNCHECK(
++      cudnnFindConvolutionForwardAlgorithm(
++         cudnnHandle,
++         inputTensorDescriptor,
++         convDescriptors->WeightsDescriptor,
++         convDescriptors->LayerDescriptor,
++         outputTensor.GetTensorDescriptor(),
++         convRequestedAlgoCount,
++         &algoCount,
++         convPerfResults
++      )
++   );
+    // we could also do it with the expert mode (cudnnFindConvolutionForwardAlgorithmEx),
+-   // but we arrive at an chicken or egg problem:
+-   // workspace size is calculated from chosen forward algorithm,
+-   // but finding a forward algorithm depends on workspace size...
+    // i.e.
+-   // Tensor_t & inputTensor = L->GetInput();
+-   // inputTensor = Tensor_t(inputTensor.GetDeviceBuffer(),{ L->GetBatchSize(), L->GetInputDepth(), L->GetInputHeight(), L->GetInputWidth() },GetTensorLayout(),0,0);
++   // create an input tensor before the inputTensorDescriptor
++   // and get the descriptor from there
++   // Tensor_t inputTensor({L->GetBatchSize(), L->GetInputDepth(), L->GetInputHeight(), L->GetInputWidth()}, MemoryLayout::RowMajor, 0, 0);
+    // CUDNNCHECK(cudnnFindConvolutionForwardAlgorithmEx(
+    //    cudnnHandle,
+-   //    inputTensorDescriptor,
++   //    inputTensor.GetTensorDescriptor(),
+    //    &inputTensor,
+    //    convDescriptors->WeightsDescriptor,
+    //    &filters,
+-   //    tempConvDescriptor,
+-   //    outputTensorDescriptor,
++   //    convDescriptors->LayerDescriptor,
++   //    outputTensor.GetTensorDescriptor(),
+    //    &outputTensor,
+    //    convRequestedAlgoCount,
+    //    &algoCount,
+    //    &convPerfResults,
+    //    &convWorkspace,
+-   //    convWorkspace->ForwardWorkspaceSize));
++   //    memLimit));  // use memLimit for workspace size
+    // instead choose either fastest or lowest memory algo as per preference
+-   int algoIdx{0};
+-   if (CNNOptions::ConvMaxWorkspaceSize != 0) {  // prefer fastest
+-      float temp_runtime{std::numeric_limits<float>::max()};
+-      for (int i = 0; i < algoCount; ++i) {
+-         if (convPerfResults[i].status != 0) continue;
+-         if (convPerfResults[i].time < temp_runtime) {
+-            temp_runtime = convPerfResults[i].time;
+-            algoIdx = i;
+-         }
+-      }
+-   } else {  // prefer smallest workspace size
+-      size_t temp_memsize{std::numeric_limits<size_t>::max()};
+-      for (int i = 0; i < algoCount; ++i) {
+-         if (convPerfResults[i].status != 0) continue;
+-         if (convPerfResults[i].memory < temp_memsize) {
+-            temp_memsize = convPerfResults[i].memory;
+-            algoIdx = i;
+-         }
+-      }
+-   }
+-   convWorkspace->AlgorithmForward = convPerfResults[algoIdx].algo;
++   convWorkspace->AlgorithmForward = convPerfResults[choose_algo(algoChoice, convPerfResults, memLimit)].algo;
+ #else
+-   // More detailed alternative: cudnnFindConvolutionForwardAlgorithm (only option in newer cuDNN versions)
+-   cudnnConvolutionFwdPreference_t preferenceFwd = (CNNOptions::ConvMaxWorkspaceSize !=0) ? CUDNN_CONVOLUTION_FWD_PREFER_FASTEST :
+-                                                   CUDNN_CONVOLUTION_FWD_NO_WORKSPACE;
+-
+-   size_t memLimit = (CNNOptions::ConvMaxWorkspaceSize > 0) ? (size_t) CNNOptions::ConvMaxWorkspaceSize : 0;
+-
+    CUDNNCHECK(cudnnGetConvolutionForwardAlgorithm(
+       cudnnHandle, inputTensorDescriptor, convDescriptors->WeightsDescriptor, convDescriptors->LayerDescriptor,
+       outputTensor.GetTensorDescriptor(), preferenceFwd,
+@@ -519,75 +557,36 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    /**
+     * I'm sure there may be a faster way, but this works
+     */
+-   convRequestedAlgoCount = 6; // reset to max number of available backward algorithms
+-   cudnnConvolutionDescriptor_t tempConvBwdDescriptor;
+-   CUDNNCHECK(cudnnCreateConvolutionDescriptor(&tempConvBwdDescriptor));
+-   cudnnTensorDescriptor_t  outputBwdTensorDescriptor;
+-   CUDNNCHECK(cudnnCreateTensorDescriptor(&outputBwdTensorDescriptor));
+-   CUDNNCHECK(cudnnSetTensor4dDescriptor(outputBwdTensorDescriptor,
+-                                             CUDNN_TENSOR_NCHW,  // Layout of the tensor in memory
+-                                             Tensor_t::GetDataType(),
+-                                             (int)L->GetBatchSize(),
+-                                             (int)L->GetInputDepth(),
+-                                             (int)L->GetInputHeight(),
+-                                             (int)L->GetInputWidth()));
+-   cudnnConvolutionBwdDataAlgoPerf_t convPerfBwdResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
++   convRequestedAlgoCount = 6;  // reset to max number of available backward algorithms
++   cudnnConvolutionBwdDataAlgoPerf_t convPerfBwdDataResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
+    CUDNNCHECK(cudnnFindConvolutionBackwardDataAlgorithm(
+       cudnnHandle,
+       convDescriptors->WeightsDescriptor,
++      activationGradients.GetTensorDescriptor(),
++      convDescriptors->LayerDescriptor,
+       activationGradientsBackwardDescriptor,
+-      tempConvBwdDescriptor,
+-      outputBwdTensorDescriptor,
+       convRequestedAlgoCount,
+       &algoCount,
+-      convPerfBwdResults));
++      convPerfBwdDataResults));
+    // we could also do it with the expert mode (cudnnFindConvolutionForwardAlgorithmEx),
+-   // but we arrive at an chicken or egg problem:
+-   // workspace size is calculated from chosen forward algorithm,
+-   // but finding a forward algorithm depends on workspace size...
+    // i.e.
+-   // Tensor_t & outputBwdTensor = L->GetInput();
+-   // outputBwdTensor = Tensor_t(outputBwdTensor.GetDeviceBuffer(),{ L->GetBatchSize(), L->GetInputDepth(), L->GetInputHeight(), L->GetInputWidth() },GetTensorLayout(),0,0);
+    // CUDNNCHECK(cudnnFindConvolutionBackwardDataAlgorithmEx(
+    //    cudnnHandle,
+    //    convDescriptors->WeightsDescriptor,
+    //    &filters,
++   //    activationGradients.GetTensorDescriptor(),
++   //    &activationGradients,
++   //    convDescriptors->LayerDescriptor,
+    //    activationGradientsBackwardDescriptor,
+-   //    &activationGradientsBackwardTensor,
+-   //    tempConvBwdDescriptor,
+-   //    outputBwdTensorDescriptor,
+-   //    &outputBwdTensor,
++   //    &inputTensor,
+    //    convRequestedAlgoCount,
+    //    &algoCount,
+    //    &convPerfBwdResults,
+    //    &convWorkspace,
+-   //    convWorkspace->ForwardWorkspaceSize));
++   //    memLimit));  // use memLimit for workspace size
+    // instead choose either fastest or lowest memory algo as per preference
+-   algoIdx = 0;
+-   if (CNNOptions::ConvMaxWorkspaceSize != 0) {  // prefer fastest
+-      float temp_runtime{std::numeric_limits<float>::max()};
+-      for (int i = 0; i < algoCount; ++i) {
+-         if (convPerfBwdResults[i].status != 0) continue;
+-         if (convPerfBwdResults[i].time < temp_runtime) {
+-            temp_runtime = convPerfBwdResults[i].time;
+-            algoIdx = i;
+-         }
+-      }
+-   } else {  // prefer smallest workspace size
+-      size_t temp_memsize{std::numeric_limits<size_t>::max()};
+-      for (int i = 0; i < algoCount; ++i) {
+-         if (convPerfBwdResults[i].status != 0) continue;
+-         if (convPerfBwdResults[i].memory < temp_memsize) {
+-            temp_memsize = convPerfBwdResults[i].memory;
+-            algoIdx = i;
+-         }
+-      }
+-   }
+-   convWorkspace->AlgorithmBackward = convPerfBwdResults[algoIdx].algo;
++   convWorkspace->AlgorithmBackward = convPerfBwdDataResults[choose_algo(algoChoice, convPerfBwdDataResults, memLimit)].algo;
+ #else
+-   cudnnConvolutionBwdDataPreference_t preferenceBwdData =
+-      (CNNOptions::ConvMaxWorkspaceSize != 0) ? CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST : CUDNN_CONVOLUTION_BWD_DATA_NO_WORKSPACE;
+-
+    CUDNNCHECK(cudnnGetConvolutionBackwardDataAlgorithm(cudnnHandle,
+                                                       convDescriptors->WeightsDescriptor,
+                                                       activationGradients.GetTensorDescriptor(),
+@@ -628,11 +627,40 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    // here should be able to use inputTensorDescriptor
+    cudnnTensorDescriptor_t activationBackwardDescriptor = inputTensorDescriptor;
+ 
+-   // cudnnConvolutionBwdFilterPreference_t preference =
+-   cudnnConvolutionBwdFilterPreference_t preferenceBwdFilter = (CNNOptions::ConvMaxWorkspaceSize != 0)
+-                                                                  ? CUDNN_CONVOLUTION_BWD_FILTER_NO_WORKSPACE
+-                                                                  : CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST;
+-
++#if (CUDNN_VERSION >= 8000)
++   /**
++    * I'm sure there may be a faster way, but this works
++    */
++   convRequestedAlgoCount = 6;  // reset to max number of available backward algorithms
++   cudnnConvolutionBwdDataAlgoPerf_t convPerfBwdFilterResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
++   CUDNNCHECK(cudnnFindConvolutionBackwardFilterAlgorithm(
++      cudnnHandle,
++      activationBackwardDescriptor,
++      activationGradients.GetTensorDescriptor(),
++      convDescriptors->LayerDescriptor,
++      convDescriptors->WeightsDescriptor,
++      convRequestedAlgoCount,
++      &algoCount,
++      convPerfBwdFilterResults));
++   // we could also do it with the expert mode (cudnnFindConvolutionForwardAlgorithmEx),
++   // i.e.
++   // CUDNNCHECK(cudnnFindConvolutionBackwardFilterAlgorithmEx(
++   //    cudnnHandle,
++   //    activationBackwardDescriptor,
++   //    &inputTensor,
++   //    activationGradients.GetTensorDescriptor(),
++   //    &activationGradients,
++   //    convDescriptors->LayerDescriptor,
++   //    convDescriptors->WeightsDescriptor,
++   //    &filters,
++   //    convRequestedAlgoCount,
++   //    &algoCount,
++   //    &convPerfBwdFilterResults,
++   //    &convWorkspace,
++   //    memLimit));  // use memLimit for workspace size
++   // instead choose either fastest or lowest memory algo as per preference
++   convWorkspace->AlgorithmBackward = convPerfBwdFilterResults[choose_algo(algoChoice, convPerfBwdFilterResults, memLimit)].algo;
++#else
+    CUDNNCHECK(cudnnGetConvolutionBackwardFilterAlgorithm(cudnnHandle,
+                                                          activationBackwardDescriptor,
+                                                          activationGradients.GetTensorDescriptor(),
+@@ -641,6 +669,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+                                                          preferenceBwdFilter,
+                                                          memLimit,
+                                                          &convWorkspace->HelperAlgorithm));
++#endif
+ 
+    std::cout << "CONV BWD Filter Algo used  is " << convWorkspace->HelperAlgorithm << std::endl;
+ 
+
+From a9d39cc9ccf9ae474d90b6671d3e0d69d4cf6872 Mon Sep 17 00:00:00 2001
+From: Konstantin Gizdov <kgizdov at gmail.com>
+Date: Wed, 22 Jul 2020 17:11:30 +0300
+Subject: [PATCH 06/10] implement correct logic behind cudnn logarithm
+ preference
+
+---
+ .../src/DNN/Architectures/Cudnn/Propagate.cu  | 20 +++++++++----------
+ 1 file changed, 10 insertions(+), 10 deletions(-)
+
+diff --git a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+index 2049e2b9195..b74c99d1a99 100644
+--- a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
++++ b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+@@ -380,18 +380,8 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+ #endif
+    // decide on algorithm preference early
+    if (CNNOptions::ConvMaxWorkspaceSize < 0) {
+-      // no workspace case
+ #if (CUDNN_VERSION >= 8000)
+-      algoChoice = no_workspace;
+-#else
+-      preferenceFwd = CUDNN_CONVOLUTION_FWD_NO_WORKSPACE;
+-      preferenceBwdData = CUDNN_CONVOLUTION_BWD_DATA_NO_WORKSPACE;
+-      preferenceBwdFilter = CUDNN_CONVOLUTION_BWD_FILTER_NO_WORKSPACE;
+-#endif
+-
+-   } else if (CNNOptions::ConvMaxWorkspaceSize == 0) {
+       // fastest overall
+-#if (CUDNN_VERSION >= 8000)
+       algoChoice = fastest;
+ #else
+       preferenceFwd = CUDNN_CONVOLUTION_FWD_PREFER_FASTEST;
+@@ -399,6 +389,16 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+       preferenceBwdFilter = CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST;
+ #endif
+ 
++   } else if (CNNOptions::ConvMaxWorkspaceSize == 0) {
++      // no workspace case
++#if (CUDNN_VERSION >= 8000)
++      algoChoice = no_workspace;
++#else
++      preferenceFwd = CUDNN_CONVOLUTION_FWD_NO_WORKSPACE;
++      preferenceBwdData = CUDNN_CONVOLUTION_BWD_DATA_NO_WORKSPACE;
++      preferenceBwdFilter = CUDNN_CONVOLUTION_BWD_FILTER_NO_WORKSPACE;
++#endif
++
+    } else {
+       // fastest in memory limit
+ #if (CUDNN_VERSION >= 8000)
+
+From 6282dfa816c7f51af5c0ecaa0065514e3f627631 Mon Sep 17 00:00:00 2001
+From: Konstantin Gizdov <kgizdov at gmail.com>
+Date: Wed, 22 Jul 2020 18:51:56 +0300
+Subject: [PATCH 07/10] use decltype instead of auto, fix typos
+
+---
+ .../src/DNN/Architectures/Cudnn/Propagate.cu  | 22 +++++++++----------
+ 1 file changed, 11 insertions(+), 11 deletions(-)
+
+diff --git a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+index b74c99d1a99..6cefd72c099 100644
+--- a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
++++ b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+@@ -343,29 +343,29 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+ #if (CUDNN_VERSION >= 8000)
+    enum algoPreference { no_workspace, fastest, workspace_limit };
+    algoPreference algoChoice;
+-   auto choose_algo = [](algoPreference const& algoPref, auto&& perfResults, size_t memLim = std::numeric_limits<size_t>::max()) -> int {
++   auto choose_algo = [](algoPreference const& algoPref, int const algoCount, decltype(perfResults) const& perfResults, size_t memLim = std::numeric_limits<size_t>::max()) -> int {
+       int algoIdx{0};
+       if (algoPref == algoPreference::fastest) {  // prefer fastest
+          float temp_runtime{std::numeric_limits<float>::max()};
+          for (int i = 0; i < algoCount; ++i) {
+-            if (PerfResults[i].status == CUDNN_STATUS_SUCCESS && PerfResults[i].time < temp_runtime) {
+-               temp_runtime = PerfResults[i].time;
++            if (perfResults[i].status == CUDNN_STATUS_SUCCESS && perfResults[i].time < temp_runtime) {
++               temp_runtime = perfResults[i].time;
+                algoIdx = i;
+             }
+          }
+       } else if (algoPref == algoPreference::workspace_limit) {  // constrain to workspace size
+          float temp_runtime{std::numeric_limits<float>::max()};
+          for (int i = 0; i < algoCount; ++i) {
+-            if (PerfResults[i].status == CUDNN_STATUS_SUCCESS && PerfResults[i].time < temp_runtime && PerfResults[i].memory <= memLim) {
+-               temp_runtime = PerfResults[i].time;
++            if (perfResults[i].status == CUDNN_STATUS_SUCCESS && perfResults[i].time < temp_runtime && perfResults[i].memory <= memLim) {
++               temp_runtime = perfResults[i].time;
+                algoIdx = i;
+             }
+          }
+       } else {  // prefer smallest workspace size
+          size_t temp_memsize{std::numeric_limits<size_t>::max()};
+          for (int i = 0; i < algoCount; ++i) {
+-            if (PerfResults[i].status == CUDNN_STATUS_SUCCESS && PerfResults[i].memory < temp_memsize) {
+-               temp_memsize = PerfResults[i].memory;
++            if (perfResults[i].status == CUDNN_STATUS_SUCCESS && perfResults[i].memory < temp_memsize) {
++               temp_memsize = perfResults[i].memory;
+                algoIdx = i;
+             }
+          }
+@@ -494,7 +494,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    //    &convWorkspace,
+    //    memLimit));  // use memLimit for workspace size
+    // instead choose either fastest or lowest memory algo as per preference
+-   convWorkspace->AlgorithmForward = convPerfResults[choose_algo(algoChoice, convPerfResults, memLimit)].algo;
++   convWorkspace->AlgorithmForward = convPerfResults[choose_algo(algoChoice, algoCount, convPerfResults, memLimit)].algo;
+ #else
+    CUDNNCHECK(cudnnGetConvolutionForwardAlgorithm(
+       cudnnHandle, inputTensorDescriptor, convDescriptors->WeightsDescriptor, convDescriptors->LayerDescriptor,
+@@ -585,7 +585,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    //    &convWorkspace,
+    //    memLimit));  // use memLimit for workspace size
+    // instead choose either fastest or lowest memory algo as per preference
+-   convWorkspace->AlgorithmBackward = convPerfBwdDataResults[choose_algo(algoChoice, convPerfBwdDataResults, memLimit)].algo;
++   convWorkspace->AlgorithmBackward = convPerfBwdDataResults[choose_algo(algoChoice, algoCount, convPerfBwdDataResults, memLimit)].algo;
+ #else
+    CUDNNCHECK(cudnnGetConvolutionBackwardDataAlgorithm(cudnnHandle,
+                                                       convDescriptors->WeightsDescriptor,
+@@ -632,7 +632,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+     * I'm sure there may be a faster way, but this works
+     */
+    convRequestedAlgoCount = 6;  // reset to max number of available backward algorithms
+-   cudnnConvolutionBwdDataAlgoPerf_t convPerfBwdFilterResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
++   cudnnConvolutionBwdFilterAlgoPerf_t convPerfBwdFilterResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
+    CUDNNCHECK(cudnnFindConvolutionBackwardFilterAlgorithm(
+       cudnnHandle,
+       activationBackwardDescriptor,
+@@ -659,7 +659,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    //    &convWorkspace,
+    //    memLimit));  // use memLimit for workspace size
+    // instead choose either fastest or lowest memory algo as per preference
+-   convWorkspace->AlgorithmBackward = convPerfBwdFilterResults[choose_algo(algoChoice, convPerfBwdFilterResults, memLimit)].algo;
++   convWorkspace->AlgorithmBackward = convPerfBwdFilterResults[choose_algo(algoChoice, algoCount, convPerfBwdFilterResults, memLimit)].algo;
+ #else
+    CUDNNCHECK(cudnnGetConvolutionBackwardFilterAlgorithm(cudnnHandle,
+                                                          activationBackwardDescriptor,
+
+From 259c1c9c4d86391d1987f6635a2aece8cae587ac Mon Sep 17 00:00:00 2001
+From: Konstantin Gizdov <kgizdov at gmail.com>
+Date: Wed, 22 Jul 2020 19:39:40 +0300
+Subject: [PATCH 08/10] assign backward filter algo to correct place
+
+---
+ tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu | 2 +-
+ 1 file changed, 1 insertion(+), 1 deletion(-)
+
+diff --git a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+index 6cefd72c099..5a80dfbc03d 100644
+--- a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
++++ b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+@@ -659,7 +659,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    //    &convWorkspace,
+    //    memLimit));  // use memLimit for workspace size
+    // instead choose either fastest or lowest memory algo as per preference
+-   convWorkspace->AlgorithmBackward = convPerfBwdFilterResults[choose_algo(algoChoice, algoCount, convPerfBwdFilterResults, memLimit)].algo;
++   convWorkspace->HelperAlgorithm = convPerfBwdFilterResults[choose_algo(algoChoice, algoCount, convPerfBwdFilterResults, memLimit)].algo;
+ #else
+    CUDNNCHECK(cudnnGetConvolutionBackwardFilterAlgorithm(cudnnHandle,
+                                                          activationBackwardDescriptor,
+
+From 2c109efea0e970b380a62f6102a286542676912a Mon Sep 17 00:00:00 2001
+From: Konstantin Gizdov <kgizdov at gmail.com>
+Date: Thu, 23 Jul 2020 17:58:58 +0300
+Subject: [PATCH 09/10] make it compile and support C++11
+
+---
+ .../src/DNN/Architectures/Cudnn/Propagate.cu  | 49 ++++++++++++-------
+ 1 file changed, 30 insertions(+), 19 deletions(-)
+
+diff --git a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+index 5a80dfbc03d..66ce64a5efc 100644
+--- a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
++++ b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+@@ -343,29 +343,37 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+ #if (CUDNN_VERSION >= 8000)
+    enum algoPreference { no_workspace, fastest, workspace_limit };
+    algoPreference algoChoice;
+-   auto choose_algo = [](algoPreference const& algoPref, int const algoCount, decltype(perfResults) const& perfResults, size_t memLim = std::numeric_limits<size_t>::max()) -> int {
++   // C++11 lambdas cannot be templated, so we have to do this HORRIBLE stuff...
++   union LocalPerf_t {
++      // these three type are absolutely equivalent
++      // and one can access them as they wish to get info
++      cudnnConvolutionFwdAlgoPerf_t * fwd;
++      cudnnConvolutionBwdFilterAlgoPerf_t * bwdFilter;
++      cudnnConvolutionBwdDataAlgoPerf_t * bwdData;
++   };
++   auto choose_algo = [](algoPreference const & algoPref, int const algoCount, LocalPerf_t const & perfResults, size_t memLim = std::numeric_limits<size_t>::max()) -> int {
+       int algoIdx{0};
+       if (algoPref == algoPreference::fastest) {  // prefer fastest
+          float temp_runtime{std::numeric_limits<float>::max()};
+          for (int i = 0; i < algoCount; ++i) {
+-            if (perfResults[i].status == CUDNN_STATUS_SUCCESS && perfResults[i].time < temp_runtime) {
+-               temp_runtime = perfResults[i].time;
++            if (perfResults.fwd[i].status == CUDNN_STATUS_SUCCESS && perfResults.fwd[i].time < temp_runtime) {
++               temp_runtime = perfResults.fwd[i].time;
+                algoIdx = i;
+             }
+          }
+       } else if (algoPref == algoPreference::workspace_limit) {  // constrain to workspace size
+          float temp_runtime{std::numeric_limits<float>::max()};
+          for (int i = 0; i < algoCount; ++i) {
+-            if (perfResults[i].status == CUDNN_STATUS_SUCCESS && perfResults[i].time < temp_runtime && perfResults[i].memory <= memLim) {
+-               temp_runtime = perfResults[i].time;
++            if (perfResults.fwd[i].status == CUDNN_STATUS_SUCCESS && perfResults.fwd[i].time < temp_runtime && perfResults.fwd[i].memory <= memLim) {
++               temp_runtime = perfResults.fwd[i].time;
+                algoIdx = i;
+             }
+          }
+       } else {  // prefer smallest workspace size
+          size_t temp_memsize{std::numeric_limits<size_t>::max()};
+          for (int i = 0; i < algoCount; ++i) {
+-            if (perfResults[i].status == CUDNN_STATUS_SUCCESS && perfResults[i].memory < temp_memsize) {
+-               temp_memsize = perfResults[i].memory;
++            if (perfResults.fwd[i].status == CUDNN_STATUS_SUCCESS && perfResults.fwd[i].memory < temp_memsize) {
++               temp_memsize = perfResults.fwd[i].memory;
+                algoIdx = i;
+             }
+          }
+@@ -461,7 +469,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    int convRequestedAlgoCount{8};  // requestedAlgoCount is setting how many algorithms to try, can be tuned, fixed for now as all available
+ 
+    int algoCount;
+-   cudnnConvolutionFwdAlgoPerf_t convPerfResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
++   cudnnConvolutionFwdAlgoPerf_t convFwdPerfResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
+    CUDNNCHECK(
+       cudnnFindConvolutionForwardAlgorithm(
+          cudnnHandle,
+@@ -471,7 +479,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+          outputTensor.GetTensorDescriptor(),
+          convRequestedAlgoCount,
+          &algoCount,
+-         convPerfResults
++         convFwdPerfResults
+       )
+    );
+    // we could also do it with the expert mode (cudnnFindConvolutionForwardAlgorithmEx),
+@@ -490,11 +498,12 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    //    &outputTensor,
+    //    convRequestedAlgoCount,
+    //    &algoCount,
+-   //    &convPerfResults,
++   //    &convFwdPerfResults,
+    //    &convWorkspace,
+    //    memLimit));  // use memLimit for workspace size
+    // instead choose either fastest or lowest memory algo as per preference
+-   convWorkspace->AlgorithmForward = convPerfResults[choose_algo(algoChoice, algoCount, convPerfResults, memLimit)].algo;
++   LocalPerf_t fwdPerfResults{convFwdPerfResults};
++   convWorkspace->AlgorithmForward = convFwdPerfResults[choose_algo(algoChoice, algoCount, fwdPerfResults, memLimit)].algo;
+ #else
+    CUDNNCHECK(cudnnGetConvolutionForwardAlgorithm(
+       cudnnHandle, inputTensorDescriptor, convDescriptors->WeightsDescriptor, convDescriptors->LayerDescriptor,
+@@ -558,7 +567,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+     * I'm sure there may be a faster way, but this works
+     */
+    convRequestedAlgoCount = 6;  // reset to max number of available backward algorithms
+-   cudnnConvolutionBwdDataAlgoPerf_t convPerfBwdDataResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
++   cudnnConvolutionBwdDataAlgoPerf_t convBwdDataPerfResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
+    CUDNNCHECK(cudnnFindConvolutionBackwardDataAlgorithm(
+       cudnnHandle,
+       convDescriptors->WeightsDescriptor,
+@@ -567,7 +576,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+       activationGradientsBackwardDescriptor,
+       convRequestedAlgoCount,
+       &algoCount,
+-      convPerfBwdDataResults));
++      convBwdDataPerfResults));
+    // we could also do it with the expert mode (cudnnFindConvolutionForwardAlgorithmEx),
+    // i.e.
+    // CUDNNCHECK(cudnnFindConvolutionBackwardDataAlgorithmEx(
+@@ -581,11 +590,12 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    //    &inputTensor,
+    //    convRequestedAlgoCount,
+    //    &algoCount,
+-   //    &convPerfBwdResults,
++   //    &convBwdDataPerfResults,
+    //    &convWorkspace,
+    //    memLimit));  // use memLimit for workspace size
+    // instead choose either fastest or lowest memory algo as per preference
+-   convWorkspace->AlgorithmBackward = convPerfBwdDataResults[choose_algo(algoChoice, algoCount, convPerfBwdDataResults, memLimit)].algo;
++   LocalPerf_t bwdDataPerfResults{convBwdDataPerfResults};
++   convWorkspace->AlgorithmBackward = convBwdDataPerfResults[choose_algo(algoChoice, algoCount, bwdDataPerfResults, memLimit)].algo;
+ #else
+    CUDNNCHECK(cudnnGetConvolutionBackwardDataAlgorithm(cudnnHandle,
+                                                       convDescriptors->WeightsDescriptor,
+@@ -632,7 +642,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+     * I'm sure there may be a faster way, but this works
+     */
+    convRequestedAlgoCount = 6;  // reset to max number of available backward algorithms
+-   cudnnConvolutionBwdFilterAlgoPerf_t convPerfBwdFilterResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
++   cudnnConvolutionBwdFilterAlgoPerf_t convBwdFilterPerfResults[convRequestedAlgoCount];  // this will store metrics to choose convolution algorithm
+    CUDNNCHECK(cudnnFindConvolutionBackwardFilterAlgorithm(
+       cudnnHandle,
+       activationBackwardDescriptor,
+@@ -641,7 +651,7 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+       convDescriptors->WeightsDescriptor,
+       convRequestedAlgoCount,
+       &algoCount,
+-      convPerfBwdFilterResults));
++      convBwdFilterPerfResults));
+    // we could also do it with the expert mode (cudnnFindConvolutionForwardAlgorithmEx),
+    // i.e.
+    // CUDNNCHECK(cudnnFindConvolutionBackwardFilterAlgorithmEx(
+@@ -655,11 +665,12 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    //    &filters,
+    //    convRequestedAlgoCount,
+    //    &algoCount,
+-   //    &convPerfBwdFilterResults,
++   //    &convBwdFilterPerfResults,
+    //    &convWorkspace,
+    //    memLimit));  // use memLimit for workspace size
+    // instead choose either fastest or lowest memory algo as per preference
+-   convWorkspace->HelperAlgorithm = convPerfBwdFilterResults[choose_algo(algoChoice, algoCount, convPerfBwdFilterResults, memLimit)].algo;
++   LocalPerf_t bwdFilterPerfResults{convBwdFilterPerfResults};
++   convWorkspace->HelperAlgorithm = convBwdFilterPerfResults[choose_algo(algoChoice, algoCount, bwdFilterPerfResults, memLimit)].algo;
+ #else
+    CUDNNCHECK(cudnnGetConvolutionBackwardFilterAlgorithm(cudnnHandle,
+                                                          activationBackwardDescriptor,
+
+From 1f1dfbbac06c29df98bdebdd9367bf566f2e7ce8 Mon Sep 17 00:00:00 2001
+From: Konstantin Gizdov <kgizdov at gmail.com>
+Date: Thu, 23 Jul 2020 21:37:33 +0300
+Subject: [PATCH 10/10] compiles completely
+
+---
+ .../src/DNN/Architectures/Cudnn/Propagate.cu  | 83 ++++++++++---------
+ 1 file changed, 46 insertions(+), 37 deletions(-)
+
+diff --git a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+index 66ce64a5efc..0694369860a 100644
+--- a/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
++++ b/tmva/tmva/src/DNN/Architectures/Cudnn/Propagate.cu
+@@ -344,41 +344,50 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    enum algoPreference { no_workspace, fastest, workspace_limit };
+    algoPreference algoChoice;
+    // C++11 lambdas cannot be templated, so we have to do this HORRIBLE stuff...
+-   union LocalPerf_t {
+-      // these three type are absolutely equivalent
+-      // and one can access them as they wish to get info
+-      cudnnConvolutionFwdAlgoPerf_t * fwd;
+-      cudnnConvolutionBwdFilterAlgoPerf_t * bwdFilter;
+-      cudnnConvolutionBwdDataAlgoPerf_t * bwdData;
+-   };
+-   auto choose_algo = [](algoPreference const & algoPref, int const algoCount, LocalPerf_t const & perfResults, size_t memLim = std::numeric_limits<size_t>::max()) -> int {
+-      int algoIdx{0};
+-      if (algoPref == algoPreference::fastest) {  // prefer fastest
+-         float temp_runtime{std::numeric_limits<float>::max()};
+-         for (int i = 0; i < algoCount; ++i) {
+-            if (perfResults.fwd[i].status == CUDNN_STATUS_SUCCESS && perfResults.fwd[i].time < temp_runtime) {
+-               temp_runtime = perfResults.fwd[i].time;
+-               algoIdx = i;
++   class LocalPerf {
++    public:
++      LocalPerf(cudnnConvolutionFwdAlgoPerf_t * fwd) {m_fwd = fwd;}
++      LocalPerf(cudnnConvolutionBwdFilterAlgoPerf_t * bwdFilter) {m_bwdFilter = bwdFilter;}
++      LocalPerf(cudnnConvolutionBwdDataAlgoPerf_t * bwdData) {m_bwdData = bwdData;}
++      size_t getMemory(int i) {return m_fwd != nullptr ? m_fwd[i].memory : m_bwdFilter != nullptr ? m_bwdFilter[i].memory : m_bwdData != nullptr ? m_bwdData[i].memory : 0;}
++      float getTime(int i) {return m_fwd != nullptr ? m_fwd[i].time : m_bwdFilter != nullptr ? m_bwdFilter[i].time : m_bwdData != nullptr ? m_bwdData[i].time : 0;}
++      cudnnStatus_t getStatus(int i) {return m_fwd != nullptr ? m_fwd[i].status : m_bwdFilter != nullptr ? m_bwdFilter[i].status : m_bwdData != nullptr ? m_bwdData[i].status : CUDNN_STATUS_BAD_PARAM;}
++      int getIdx(algoPreference const & algoPref, int const algoCount, size_t memLim = std::numeric_limits<size_t>::max()) {
++         int algoIdx{0};
++         if (algoPref == algoPreference::fastest) {  // prefer fastest
++            float temp_runtime{std::numeric_limits<float>::max()};
++            for (int i = 0; i < algoCount; ++i) {
++               if (getStatus(i) == CUDNN_STATUS_SUCCESS && getTime(i) < temp_runtime) {
++                  temp_runtime = getTime(i);
++                  algoIdx = i;
++               }
+             }
+-         }
+-      } else if (algoPref == algoPreference::workspace_limit) {  // constrain to workspace size
+-         float temp_runtime{std::numeric_limits<float>::max()};
+-         for (int i = 0; i < algoCount; ++i) {
+-            if (perfResults.fwd[i].status == CUDNN_STATUS_SUCCESS && perfResults.fwd[i].time < temp_runtime && perfResults.fwd[i].memory <= memLim) {
+-               temp_runtime = perfResults.fwd[i].time;
+-               algoIdx = i;
++         } else if (algoPref == algoPreference::workspace_limit) {  // constrain to workspace size
++            float temp_runtime{std::numeric_limits<float>::max()};
++            for (int i = 0; i < algoCount; ++i) {
++               if (getStatus(i) == CUDNN_STATUS_SUCCESS && getTime(i) < temp_runtime && getMemory(i) <= memLim) {
++                  temp_runtime = getTime(i);
++                  algoIdx = i;
++               }
+             }
+-         }
+-      } else {  // prefer smallest workspace size
+-         size_t temp_memsize{std::numeric_limits<size_t>::max()};
+-         for (int i = 0; i < algoCount; ++i) {
+-            if (perfResults.fwd[i].status == CUDNN_STATUS_SUCCESS && perfResults.fwd[i].memory < temp_memsize) {
+-               temp_memsize = perfResults.fwd[i].memory;
+-               algoIdx = i;
++         } else {  // prefer smallest workspace size
++            size_t temp_memsize{std::numeric_limits<size_t>::max()};
++            for (int i = 0; i < algoCount; ++i) {
++               if (getStatus(i) == CUDNN_STATUS_SUCCESS && getMemory(i) < temp_memsize) {
++                  temp_memsize = getMemory(i);
++                  algoIdx = i;
++               }
+             }
+          }
+-      }
+-      return algoIdx;
++         return algoIdx;
++      };
++    private:
++      LocalPerf();
++      // these three type are absolutely equivalent
++      // and one can access them as they wish to get info
++      cudnnConvolutionFwdAlgoPerf_t * m_fwd;
++      cudnnConvolutionBwdFilterAlgoPerf_t * m_bwdFilter;
++      cudnnConvolutionBwdDataAlgoPerf_t * m_bwdData;
+    };
+ #else
+    // More detailed alternative: cudnnFindConvolutionForwardAlgorithm (only option in newer cuDNN versions)
+@@ -502,8 +511,8 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    //    &convWorkspace,
+    //    memLimit));  // use memLimit for workspace size
+    // instead choose either fastest or lowest memory algo as per preference
+-   LocalPerf_t fwdPerfResults{convFwdPerfResults};
+-   convWorkspace->AlgorithmForward = convFwdPerfResults[choose_algo(algoChoice, algoCount, fwdPerfResults, memLimit)].algo;
++   LocalPerf fwdPerfResults{convFwdPerfResults};
++   convWorkspace->AlgorithmForward = convFwdPerfResults[fwdPerfResults.getIdx(algoChoice, algoCount, memLimit)].algo;
+ #else
+    CUDNNCHECK(cudnnGetConvolutionForwardAlgorithm(
+       cudnnHandle, inputTensorDescriptor, convDescriptors->WeightsDescriptor, convDescriptors->LayerDescriptor,
+@@ -594,8 +603,8 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    //    &convWorkspace,
+    //    memLimit));  // use memLimit for workspace size
+    // instead choose either fastest or lowest memory algo as per preference
+-   LocalPerf_t bwdDataPerfResults{convBwdDataPerfResults};
+-   convWorkspace->AlgorithmBackward = convBwdDataPerfResults[choose_algo(algoChoice, algoCount, bwdDataPerfResults, memLimit)].algo;
++   LocalPerf bwdDataPerfResults{convBwdDataPerfResults};
++   convWorkspace->AlgorithmBackward = convBwdDataPerfResults[bwdDataPerfResults.getIdx(algoChoice, algoCount, memLimit)].algo;
+ #else
+    CUDNNCHECK(cudnnGetConvolutionBackwardDataAlgorithm(cudnnHandle,
+                                                       convDescriptors->WeightsDescriptor,
+@@ -669,8 +678,8 @@ void TCudnn<AFloat>::InitializeConvWorkspace(TWorkspace * & workspace,
+    //    &convWorkspace,
+    //    memLimit));  // use memLimit for workspace size
+    // instead choose either fastest or lowest memory algo as per preference
+-   LocalPerf_t bwdFilterPerfResults{convBwdFilterPerfResults};
+-   convWorkspace->HelperAlgorithm = convBwdFilterPerfResults[choose_algo(algoChoice, algoCount, bwdFilterPerfResults, memLimit)].algo;
++   LocalPerf bwdFilterPerfResults{convBwdFilterPerfResults};
++   convWorkspace->HelperAlgorithm = convBwdFilterPerfResults[bwdFilterPerfResults.getIdx(algoChoice, algoCount, memLimit)].algo;
+ #else
+    CUDNNCHECK(cudnnGetConvolutionBackwardFilterAlgorithm(cudnnHandle,
+                                                          activationBackwardDescriptor,

Deleted: add_missing_include_ROOT-10315.patch
===================================================================
--- add_missing_include_ROOT-10315.patch	2020-07-23 20:13:52 UTC (rev 665089)
+++ add_missing_include_ROOT-10315.patch	2020-07-23 20:24:09 UTC (rev 665090)
@@ -1,22 +0,0 @@
-From e31986c6523253f3160753b202e22a1c321e2b1a Mon Sep 17 00:00:00 2001
-From: Sergey Linev <S.Linev at gsi.de>
-Date: Thu, 12 Sep 2019 14:06:13 +0200
-Subject: [PATCH] [rbrowser] add missing include ROOT-10315
-
----
- gui/browserv7/inc/ROOT/RBrowserItem.hxx | 3 +++
- 1 file changed, 3 insertions(+)
-
-diff --git a/gui/browserv7/inc/ROOT/RBrowserItem.hxx b/gui/browserv7/inc/ROOT/RBrowserItem.hxx
-index b21bc87b4bf..e9ae97e90c0 100644
---- a/gui/browserv7/inc/ROOT/RBrowserItem.hxx
-+++ b/gui/browserv7/inc/ROOT/RBrowserItem.hxx
-@@ -17,6 +17,9 @@
- #ifndef ROOT7_RBrowserItem
- #define ROOT7_RBrowserItem
- 
-+#include <string>
-+#include <vector>
-+
- namespace ROOT {
- namespace Experimental {

Deleted: fix_relative_paths.patch
===================================================================
--- fix_relative_paths.patch	2020-07-23 20:13:52 UTC (rev 665089)
+++ fix_relative_paths.patch	2020-07-23 20:24:09 UTC (rev 665090)
@@ -1,39 +0,0 @@
-diff --git a/cmake/modules/RootConfiguration.cmake b/cmake/modules/RootConfiguration.cmake
-index 41a9cffc78..1c70ce8e39 100644
---- a/cmake/modules/RootConfiguration.cmake
-+++ b/cmake/modules/RootConfiguration.cmake
-@@ -678,22 +678,30 @@ file(RELATIVE_PATH ROOT_CMAKE_TO_INCLUDE_DIR "${CMAKE_INSTALL_FULL_CMAKEDIR}" "$
- file(RELATIVE_PATH ROOT_CMAKE_TO_LIB_DIR "${CMAKE_INSTALL_FULL_CMAKEDIR}" "${CMAKE_INSTALL_FULL_LIBDIR}")
- file(RELATIVE_PATH ROOT_CMAKE_TO_BIN_DIR "${CMAKE_INSTALL_FULL_CMAKEDIR}" "${CMAKE_INSTALL_FULL_BINDIR}")
- 
-+# '_' prefixed variables are used to construct the paths,
-+# while the normal variants evaluate to full paths at runtime
- set(ROOT_INCLUDE_DIR_SETUP "
- # ROOT configured for the install with relative paths, so use these
--get_filename_component(ROOT_INCLUDE_DIRS \"\${_thisdir}/${ROOT_CMAKE_TO_INCLUDE_DIR}\" ABSOLUTE)
-+get_filename_component(_ROOT_INCLUDE_DIRS \"\${_thisdir}/${ROOT_CMAKE_TO_INCLUDE_DIR}\" REALPATH)
-+# resolve relative paths to absolute system paths
-+get_filename_component(ROOT_INCLUDE_DIRS \"\${_ROOT_INCLUDE_DIRS}\" REALPATH)
- ")
- set(ROOT_LIBRARY_DIR_SETUP "
- # ROOT configured for the install with relative paths, so use these
--get_filename_component(ROOT_LIBRARY_DIR \"\${_thisdir}/${ROOT_CMAKE_TO_LIB_DIR}\" ABSOLUTE)
-+get_filename_component(_ROOT_LIBRARY_DIR \"\${_thisdir}/${ROOT_CMAKE_TO_LIB_DIR}\" REALPATH)
-+# resolve relative paths to absolute system paths
-+get_filename_component(ROOT_LIBRARY_DIR \"\${_ROOT_LIBRARY_DIR}\" REALPATH)
- ")
- set(ROOT_BINDIR_SETUP "
- # ROOT configured for the install with relative paths, so use these
--get_filename_component(ROOT_BINDIR \"\${_thisdir}/${ROOT_CMAKE_TO_BIN_DIR}\" ABSOLUTE)
-+get_filename_component(_ROOT_BINDIR \"\${_thisdir}/${ROOT_CMAKE_TO_BIN_DIR}\" REALPATH)
-+# resolve relative paths to absolute system paths
-+get_filename_component(ROOT_BINDIR \"\${_ROOT_BINDIR}\" REALPATH)
- ")
- # Deprecated value ROOT_BINARY_DIR
- set(ROOT_BINARY_DIR_SETUP "
- # Deprecated value, please don't use it and use ROOT_BINDIR instead.
--get_filename_component(ROOT_BINARY_DIR \"\${ROOT_BINDIR}\" ABSOLUTE)
-+get_filename_component(ROOT_BINARY_DIR \"\${ROOT_BINDIR}\" REALPATH)
- ")
- 
- # used by ROOTConfig.cmake from the build directory

Modified: nbman-for-arch.patch
===================================================================
--- nbman-for-arch.patch	2020-07-23 20:13:52 UTC (rev 665089)
+++ nbman-for-arch.patch	2020-07-23 20:24:09 UTC (rev 665090)
@@ -1,6 +1,6 @@
-diff -aur root-6.20.04-old/main/src/nbmain.cxx root-6.20.04-new/main/src/nbmain.cxx
---- root-6.20.04-old/main/src/nbmain.cxx	2020-04-14 00:54:43.902408710 +0300
-+++ root-6.20.04-new/main/src/nbmain.cxx	2020-04-14 01:00:48.022396130 +0300
+diff --color -aur root-6.22.00-old/main/src/nbmain.cxx root-6.22.00-new/main/src/nbmain.cxx
+--- root-6.22.00-old/main/src/nbmain.cxx	2020-07-20 15:26:53.983725609 +0300
++++ root-6.22.00-new/main/src/nbmain.cxx	2020-07-20 15:29:53.940386060 +0300
 @@ -33,10 +33,6 @@
  #define NB_OPT             "notebook"
  #define JUPYTER_CONF_DIR_V "JUPYTER_CONFIG_DIR"
@@ -9,10 +9,10 @@
 -#define ROOTNB_DIR         ".rootnb"
 -#define COMMIT_FILE        ".rootcommit"
 -#define JUPYTER_CONFIG     "jupyter_notebook_config.py"
- 
+
  using namespace std;
- 
-@@ -49,160 +45,12 @@
+
+@@ -46,161 +46,12 @@
  #endif
  
  ////////////////////////////////////////////////////////////////////////////////
@@ -60,10 +60,11 @@
 -
 -   // Copy files in source to dest
 -   TSystemDirectory dir(source.c_str(), source.c_str());
--   TList *files = dir.GetListOfFiles();
+-   std::unique_ptr<TList> files;
+-   files.reset(dir.GetListOfFiles());
 -   if (files) {
 -      TSystemFile *file;
--      TListIter it(files);
+-      TListIter it(files.get());
 -      while ((file = (TSystemFile*)it())) {
 -         TString s = file->GetName();
 -         string fname(s.Data());


More information about the arch-commits mailing list