Skip to content

Commit

Permalink
[HIPIFY][ROCm#584][DNN][MIOpen][refactor] cuDNN -> MIOpen - Part 8
Browse files Browse the repository at this point in the history 
+ [refactor] Introduced `ArgCastStruct` with additional fields `isToRoc` and `isToMIOpen` (both are `false` by default) for the correct argument casting in `roc*` and `miopen*` functions
+ Removed the penultimate (5th) arg `cudnnConvolutionFwdAlgo_t algo` from the `miopenConvolutionForwardGetWorkSpaceSize` function call after hipification of `cudnnGetConvolutionForwardWorkspaceSize` due to the absence of that argument in `miopenConvolutionForwardGetWorkSpaceSize`; no warning is emitted
+ Updated the synthetic test `cudnn2miopen.cu` accordingly
  • Loading branch information
@emankov
emankov committed Apr 1, 2023
1 parent b74240c commit eed2258
Show file tree
Hide file tree
Showing 4 changed files with 132 additions and 25 deletions.
4 changes: 2 additions & 2 deletions src/CUDA2HIP_Perl.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand Down
8 changes: 7 additions & 1 deletion src/CUDA2HIP_Scripting.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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 {

Expand Down
140 changes: 121 additions & 19 deletions src/HipifyAction.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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";
Expand All @@ -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) {
Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -662,7 +763,8 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
sCudaGraphExecMemcpyNodeSetParamsFromSymbol,
sCuOccupancyMaxPotentialBlockSize,
sCuOccupancyMaxPotentialBlockSizeWithFlags,
sCudaGetTextureReference
sCudaGetTextureReference,
sCudnnGetConvolutionForwardWorkspaceSize
)
)
)
Expand Down
5 changes: 2 additions & 3 deletions tests/unit_tests/synthetic/libraries/cudnn2miopen.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down

0 comments on commit eed2258

Please sign in to comment.