From 3483f9bd0d053f0330a7f15d5f56d5063ecc3f04 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 4 Apr 2023 12:13:21 +0200 Subject: [PATCH] [HIPIFY][#584][DNN][MIOpen] cuDNN -> MIOpen - Part 12 - cuDNN LRN functions + Continued supporting hipification to MIOpen based on `miopen.h` + Updated the synthetic test `cudnn2miopen.cu` accordingly + [IMP] `cudnnLRNCrossChannelForward` and `cudnnLRNCrossChannelBackward` do not have a correspondence to `miopenLRNForward` and `miopenLRNBackward` - to discuss with the team [ToDo] + [feature] Add a new type of transformation with declaring a var before the function call to add that var reference as an arg to the below function call + Add a referrence to miopenLRNMode_t as a 2nd arg for `miopenGetLRNDescriptor` function call --- src/CUDA2HIP_DNN_API_functions.cpp | 8 +++--- src/HipifyAction.cpp | 13 ++++++++- .../synthetic/libraries/cudnn2miopen.cu | 27 +++++++++++++++++++ 3 files changed, 43 insertions(+), 5 deletions(-) diff --git a/src/CUDA2HIP_DNN_API_functions.cpp b/src/CUDA2HIP_DNN_API_functions.cpp index b5a40ae3..3868a7d3 100644 --- a/src/CUDA2HIP_DNN_API_functions.cpp +++ b/src/CUDA2HIP_DNN_API_functions.cpp @@ -160,10 +160,10 @@ const std::map CUDA_DNN_FUNCTION_MAP { {"cudnnGetActivationDescriptorSwishBeta", {"hipdnnGetActivationDescriptorSwishBeta", "", CONV_LIB_FUNC, API_DNN, 2, HIP_UNSUPPORTED}}, // cuDNN LRN functions - {"cudnnCreateLRNDescriptor", {"hipdnnCreateLRNDescriptor", "", CONV_LIB_FUNC, API_DNN, 2}}, - {"cudnnSetLRNDescriptor", {"hipdnnSetLRNDescriptor", "", CONV_LIB_FUNC, API_DNN, 2}}, - {"cudnnGetLRNDescriptor", {"hipdnnGetLRNDescriptor", "", CONV_LIB_FUNC, API_DNN, 2}}, - {"cudnnDestroyLRNDescriptor", {"hipdnnDestroyLRNDescriptor", "", CONV_LIB_FUNC, API_DNN, 2}}, + {"cudnnCreateLRNDescriptor", {"hipdnnCreateLRNDescriptor", "miopenCreateLRNDescriptor", CONV_LIB_FUNC, API_DNN, 2}}, + {"cudnnSetLRNDescriptor", {"hipdnnSetLRNDescriptor", "miopenSetLRNDescriptor", CONV_LIB_FUNC, API_DNN, 2}}, + {"cudnnGetLRNDescriptor", {"hipdnnGetLRNDescriptor", "miopenGetLRNDescriptor", CONV_LIB_FUNC, API_DNN, 2}}, + {"cudnnDestroyLRNDescriptor", {"hipdnnDestroyLRNDescriptor", "miopenDestroyLRNDescriptor", CONV_LIB_FUNC, API_DNN, 2}}, {"cudnnLRNCrossChannelForward", {"hipdnnLRNCrossChannelForward", "", CONV_LIB_FUNC, API_DNN, 2}}, {"cudnnLRNCrossChannelBackward", {"hipdnnLRNCrossChannelBackward", "", CONV_LIB_FUNC, API_DNN, 2}}, diff --git a/src/HipifyAction.cpp b/src/HipifyAction.cpp index 6139c192..d5abe57b 100644 --- a/src/HipifyAction.cpp +++ b/src/HipifyAction.cpp @@ -74,6 +74,7 @@ const std::string sCudnnSetPooling2dDescriptor = "cudnnSetPooling2dDescriptor"; const std::string sCudnnGetPooling2dDescriptor = "cudnnGetPooling2dDescriptor"; const std::string sCudnnSetPoolingNdDescriptor = "cudnnSetPoolingNdDescriptor"; const std::string sCudnnGetPoolingNdDescriptor = "cudnnGetPoolingNdDescriptor"; +const std::string sCudnnSetLRNDescriptor = "cudnnSetLRNDescriptor"; // Matchers' names const StringRef sCudaLaunchKernel = "cudaLaunchKernel"; const StringRef sCudaHostFuncCall = "cudaHostFuncCall"; @@ -264,6 +265,15 @@ std::map FuncArgCasts { true } }, + {sCudnnSetLRNDescriptor, + { + { + {1, {e_add_const_argument, cw_None, "miopenLRNCrossChannel"}} + }, + true, + true + } + }, }; void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) { @@ -846,7 +856,8 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi sCudnnSetPooling2dDescriptor, sCudnnGetPooling2dDescriptor, sCudnnSetPoolingNdDescriptor, - sCudnnGetPoolingNdDescriptor + sCudnnGetPoolingNdDescriptor, + sCudnnSetLRNDescriptor ) ) ) diff --git a/tests/unit_tests/synthetic/libraries/cudnn2miopen.cu b/tests/unit_tests/synthetic/libraries/cudnn2miopen.cu index b049b090..6ff27412 100644 --- a/tests/unit_tests/synthetic/libraries/cudnn2miopen.cu +++ b/tests/unit_tests/synthetic/libraries/cudnn2miopen.cu @@ -436,5 +436,32 @@ int main() { // CHECK: status = miopenDestroyPoolingDescriptor(poolingDescriptor); status = cudnnDestroyPoolingDescriptor(poolingDescriptor); + unsigned lrnN = 0; + double lrnAlpha = 0.0f; + double lrnBeta = 0.0f; + double lrnK = 0.0f; + + // CUDA: cudnnStatus_t CUDNNWINAPI cudnnCreateLRNDescriptor(cudnnLRNDescriptor_t* normDesc); + // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenCreateLRNDescriptor(miopenLRNDescriptor_t* lrnDesc); + // CHECK: status = miopenCreateLRNDescriptor(&LRNDescriptor); + status = cudnnCreateLRNDescriptor(&LRNDescriptor); + + // CUDA: cudnnStatus_t CUDNNWINAPI cudnnSetLRNDescriptor(cudnnLRNDescriptor_t normDesc, unsigned lrnN, double lrnAlpha, double lrnBeta, double lrnK); + // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenSetLRNDescriptor(const miopenLRNDescriptor_t lrnDesc, miopenLRNMode_t mode, unsigned int lrnN, double lrnAlpha, double lrnBeta, double lrnK); + // CHECK: status = miopenSetLRNDescriptor(LRNDescriptor, miopenLRNCrossChannel, lrnN, lrnAlpha, lrnBeta, lrnK); + status = cudnnSetLRNDescriptor(LRNDescriptor, lrnN, lrnAlpha, lrnBeta, lrnK); + + // TODO: add a referrence to miopenLRNMode_t as a 2nd arg + // TODO: [feature] Add a new type of transformation with declaring a var before the function call to add that var referrence as an arg to the below function call + // CUDA: cudnnStatus_t CUDNNWINAPI cudnnGetLRNDescriptor(cudnnLRNDescriptor_t normDesc, unsigned* lrnN, double* lrnAlpha, double* lrnBeta, double* lrnK); + // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenGetLRNDescriptor(const miopenLRNDescriptor_t lrnDesc, miopenLRNMode_t* mode, unsigned int* lrnN, double* lrnAlpha, double* lrnBeta, double* lrnK); + // CHECK: status = miopenGetLRNDescriptor(LRNDescriptor, &lrnN, &lrnAlpha, &lrnBeta, &lrnK); + status = cudnnGetLRNDescriptor(LRNDescriptor, &lrnN, &lrnAlpha, &lrnBeta, &lrnK); + + // CUDA: cudnnStatus_t CUDNNWINAPI cudnnDestroyLRNDescriptor(cudnnLRNDescriptor_t lrnDesc); + // MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenDestroyLRNDescriptor(miopenLRNDescriptor_t lrnDesc); + // CHECK: status = miopenDestroyLRNDescriptor(LRNDescriptor); + status = cudnnDestroyLRNDescriptor(LRNDescriptor); + return 0; }