Merge pull request #821 from emankov/HIPIFY

[HIPIFY][#584][DNN][MIOpen][refactor] cuDNN -> MIOpen - Part 8
ROCm · Apr 1, 2023 · 31d10de · 31d10de
2 parents 05edfbe + eed2258
commit 31d10de
Show file tree

Hide file tree

Showing 4 changed files with 132 additions and 25 deletions.
diff --git a/src/CUDA2HIP_Perl.cpp b/src/CUDA2HIP_Perl.cpp
@@ -487,8 +487,8 @@ namespace perl {
     set<string> ReinterpretFunctions1;
     set<string> RemoveArgFunctions3;
     for (auto f : FuncArgCasts) {
-      auto casts = f.second;
-      for (auto c : casts) {
+      auto castStruct = f.second;
+      for (auto c : castStruct.castMap) {
         switch (c.first) {
           case 0:
             switch (c.second.castType) {

diff --git a/src/CUDA2HIP_Scripting.h b/src/CUDA2HIP_Scripting.h
@@ -43,10 +43,16 @@ namespace hipify {
   };
 
   typedef std::map<unsigned, CastInfo> ArgCastMap;
+
+  struct ArgCastStruct {
+    ArgCastMap castMap;
+    bool isToRoc = false;
+    bool isToMIOpen = false;
+  };
 }
 
 extern std::string getCastType(hipify::CastTypes c);
-extern std::map<std::string, hipify::ArgCastMap> FuncArgCasts;
+extern std::map<std::string, hipify::ArgCastStruct> FuncArgCasts;
 
 namespace perl {
 

diff --git a/src/HipifyAction.cpp b/src/HipifyAction.cpp
@@ -67,6 +67,7 @@ const std::string sCudaGraphExecMemcpyNodeSetParamsFromSymbol = "cudaGraphExecMe
 const std::string sCuOccupancyMaxPotentialBlockSize = "cuOccupancyMaxPotentialBlockSize";
 const std::string sCuOccupancyMaxPotentialBlockSizeWithFlags = "cuOccupancyMaxPotentialBlockSizeWithFlags";
 const std::string sCudaGetTextureReference = "cudaGetTextureReference";
+const std::string sCudnnGetConvolutionForwardWorkspaceSize = "cudnnGetConvolutionForwardWorkspaceSize";
 // Matchers' names
 const StringRef sCudaLaunchKernel = "cudaLaunchKernel";
 const StringRef sCudaHostFuncCall = "cudaHostFuncCall";
@@ -87,22 +88,121 @@ std::string getCastType(hipify::CastTypes c) {
   }
 }
 
-std::map<std::string, ArgCastMap> FuncArgCasts {
-  {sCudaMemcpyToSymbol, {{0, {e_HIP_SYMBOL, cw_None}}}},
-  {sCudaMemcpyToSymbolAsync, {{0, {e_HIP_SYMBOL, cw_None}}}},
-  {sCudaGetSymbolSize, {{1, {e_HIP_SYMBOL, cw_None}}}},
-  {sCudaGetSymbolAddress, {{1, {e_HIP_SYMBOL, cw_None}}}},
-  {sCudaMemcpyFromSymbol, {{1, {e_HIP_SYMBOL, cw_None}}}},
-  {sCudaMemcpyFromSymbolAsync, {{1, {e_HIP_SYMBOL, cw_None}}}},
-  {sCudaGraphAddMemcpyNodeToSymbol, {{4, {e_HIP_SYMBOL, cw_None}}}},
-  {sCudaGraphAddMemcpyNodeFromSymbol, {{5, {e_HIP_SYMBOL, cw_None}}}},
-  {sCudaGraphMemcpyNodeSetParamsToSymbol, {{1, {e_HIP_SYMBOL, cw_None}}}},
-  {sCudaGraphMemcpyNodeSetParamsFromSymbol, {{2, {e_HIP_SYMBOL, cw_None}}}},
-  {sCudaGraphExecMemcpyNodeSetParamsToSymbol, {{2, {e_HIP_SYMBOL, cw_None}}}},
-  {sCudaGraphExecMemcpyNodeSetParamsFromSymbol, {{3, {e_HIP_SYMBOL, cw_None}}}},
-  {sCudaGetTextureReference, {{1, {e_HIP_SYMBOL, cw_None}}}},
-  {sCuOccupancyMaxPotentialBlockSize, {{3, {e_remove_argument, cw_DataLoss}}}},
-  {sCuOccupancyMaxPotentialBlockSizeWithFlags, {{3, {e_remove_argument, cw_DataLoss}}}},
+std::map<std::string, ArgCastStruct> FuncArgCasts {
+  {sCudaMemcpyToSymbol,
+    {
+      {
+        {0, {e_HIP_SYMBOL, cw_None}}
+      }
+    }
+  },
+  {sCudaMemcpyToSymbolAsync,
+    {
+      {
+        {0, {e_HIP_SYMBOL, cw_None}}
+      }
+    }
+  },
+  {sCudaGetSymbolSize,
+    {
+      {
+        {1, {e_HIP_SYMBOL, cw_None}}
+      }
+    }
+  },
+  {sCudaGetSymbolAddress,
+    {
+      {
+        {1, {e_HIP_SYMBOL, cw_None}}
+      }
+    }
+  },
+  {sCudaMemcpyFromSymbol,
+    {
+      {
+        {1, {e_HIP_SYMBOL, cw_None}}
+      }
+    }
+  },
+  {sCudaMemcpyFromSymbolAsync,
+    {
+      {
+        {1, {e_HIP_SYMBOL, cw_None}}
+      }
+    }
+  },
+  {sCudaGraphAddMemcpyNodeToSymbol,
+    {
+      {
+        {4, {e_HIP_SYMBOL, cw_None}}
+      }
+    }
+  },
+  {sCudaGraphAddMemcpyNodeFromSymbol,
+    {
+      {
+        {5, {e_HIP_SYMBOL, cw_None}}
+      }
+    }
+  },
+  {sCudaGraphMemcpyNodeSetParamsToSymbol,
+    {
+      {
+        {1, {e_HIP_SYMBOL, cw_None}}
+      }
+    }
+  },
+  {sCudaGraphMemcpyNodeSetParamsFromSymbol,
+    {
+      {
+        {2, {e_HIP_SYMBOL, cw_None}}
+      }
+    }
+  },
+  {sCudaGraphExecMemcpyNodeSetParamsToSymbol,
+    {
+      {
+        {2, {e_HIP_SYMBOL, cw_None}}
+      }
+    }
+  },
+  {sCudaGraphExecMemcpyNodeSetParamsFromSymbol,
+    {
+      {
+        {3, {e_HIP_SYMBOL, cw_None}}
+      }
+    }
+  },
+  {sCudaGetTextureReference,
+    {
+      {
+        {1, {e_HIP_SYMBOL, cw_None}}
+      }
+    }
+  },
+  {sCuOccupancyMaxPotentialBlockSize,
+    {
+      {
+        {3, {e_remove_argument, cw_DataLoss}}
+      }
+    }
+  },
+  {sCuOccupancyMaxPotentialBlockSizeWithFlags,
+    {
+      {
+        {3, {e_remove_argument, cw_DataLoss}}
+      }
+    }
+  },
+  {sCudnnGetConvolutionForwardWorkspaceSize,
+    {
+      {
+        {5, {e_remove_argument, cw_None}}
+      },
+      true,
+      true
+    }
+  },
 };
 
 void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
@@ -539,9 +639,10 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result)
     std::string sName = funcDcl->getDeclName().getAsString();
     auto it = FuncArgCasts.find(sName);
     if (it == FuncArgCasts.end()) return false;
+    auto castStruct = it->second;
+    if (castStruct.isToMIOpen != TranslateToMIOpen || castStruct.isToRoc != TranslateToRoc) return false;
     clang::LangOptions DefaultLangOptions;
-    auto casts = it->second;
-    for (auto c : casts) {
+    for (auto c : castStruct.castMap) {
       unsigned int argNum = c.first;
       clang::SmallString<40> XStr;
       llvm::raw_svector_ostream OS(XStr);
@@ -662,7 +763,8 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
             sCudaGraphExecMemcpyNodeSetParamsFromSymbol,
             sCuOccupancyMaxPotentialBlockSize,
             sCuOccupancyMaxPotentialBlockSizeWithFlags,
-            sCudaGetTextureReference
+            sCudaGetTextureReference,
+            sCudnnGetConvolutionForwardWorkspaceSize
           )
         )
       )

diff --git a/tests/unit_tests/synthetic/libraries/cudnn2miopen.cu b/tests/unit_tests/synthetic/libraries/cudnn2miopen.cu
@@ -335,12 +335,11 @@ int main() {
   // CHECK: status = miopenFindConvolutionForwardAlgorithm(handle, xD, x, filterDescriptor, W, convolutionDescriptor, yD, y, requestedAlgoCount, &returnedAlgoCount, &ConvolutionFwdAlgoPerf_t, workSpace, workSpaceSizeInBytes);
   status = cudnnFindConvolutionForwardAlgorithmEx(handle, xD, x, filterDescriptor, W, convolutionDescriptor, yD, y, requestedAlgoCount, &returnedAlgoCount, &ConvolutionFwdAlgoPerf_t, workSpace, workSpaceSizeInBytes);
 
-  // TODO: remove the penultimate arg (cudnnConvolutionFwdAlgo_t algo)
   // TODO: swap 2 and 3 arguments (const miopenTensorDescriptor_t wDesc and const miopenTensorDescriptor_t xDesc)
   // CUDA: cudnnStatus_t CUDNNWINAPI cudnnGetConvolutionForwardWorkspaceSize(cudnnHandle_t handle, const cudnnTensorDescriptor_t xDesc, const cudnnFilterDescriptor_t wDesc, const cudnnConvolutionDescriptor_t convDesc, const cudnnTensorDescriptor_t yDesc, cudnnConvolutionFwdAlgo_t algo, size_t* sizeInBytes);
   // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenConvolutionForwardGetWorkSpaceSize(miopenHandle_t handle, const miopenTensorDescriptor_t wDesc, const miopenTensorDescriptor_t xDesc, const miopenConvolutionDescriptor_t convDesc, const miopenTensorDescriptor_t yDesc, size_t* workSpaceSize);
-  // CHECK: status = miopenConvolutionForwardGetWorkSpaceSize(handle, xD, filterDescriptor, convolutionDescriptor, yD, convolutionFwdAlgo , &workSpaceSizeInBytes);
-  status = cudnnGetConvolutionForwardWorkspaceSize(handle, xD, filterDescriptor, convolutionDescriptor, yD, convolutionFwdAlgo , &workSpaceSizeInBytes);
+  // CHECK: status = miopenConvolutionForwardGetWorkSpaceSize(handle, xD, filterDescriptor, convolutionDescriptor, yD,  &workSpaceSizeInBytes);
+  status = cudnnGetConvolutionForwardWorkspaceSize(handle, xD, filterDescriptor, convolutionDescriptor, yD, convolutionFwdAlgo, &workSpaceSizeInBytes);
 
   // TODO: swap correstly last 5 arguments
   // CUDA: cudnnStatus_t CUDNNWINAPI cudnnConvolutionForward(cudnnHandle_t handle, const void* alpha, const cudnnTensorDescriptor_t xDesc, const void* x, const cudnnFilterDescriptor_t wDesc, const void* w, const cudnnConvolutionDescriptor_t convDesc, cudnnConvolutionFwdAlgo_t algo, void* workSpace, size_t workSpaceSizeInBytes, const void* beta, const cudnnTensorDescriptor_t yDesc, void* y);