Skip to content

[HIPIFY][#837][#2062][RFC][feature] Implement e_insert_new_argument for APIs need variable declarations - Part 1 #2415

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
emankov merged 3 commits into from
Mar 18, 2026
+93 −8
Merged

[HIPIFY][#837][#2062][RFC][feature] Implement e_insert_new_argument for APIs need variable declarations - Part 1 #2415

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
58926dd
[HIPIFY][#837][#2062][feature][RFC] Implement e_insert_new_argument f…
ranapratap55 Feb 25, 2026
873b2b2
[HIPIFY][#837][#2062][feature] Fix scoping, add variable initializati…
ranapratap55 Mar 6, 2026
750aeb7
[HIPIFY] Add hipify_ prefix for generated variable names in e_insert_…
ranapratap55 Mar 18, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 7 additions & 0 deletions src/CUDA2HIP_Scripting.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,6 +41,7 @@ namespace hipify {
e_add_var_argument,
e_move_argument,
e_replace_argument_with_const,
e_insert_new_argument,
};

enum OverloadTypes {
Expand All @@ -50,6 +51,7 @@ namespace hipify {
enum CastWarning {
cw_None,
cw_DataLoss,
cw_NeedsNewArgDecl,
};

enum OverloadWarning {
Expand All @@ -59,9 +61,14 @@ namespace hipify {
struct CastInfo {
CastTypes castType = e_HIP_SYMBOL;
CastWarning castWarn = cw_None;
// For e_add_const_argument / e_replace_argument_with_const: the literal value to insert or replace.
// For e_insert_new_argument: the variable name to declare and pass as an argument.
std::string constValToAddOrReplace = "";
Comment thread
ranapratap55 marked this conversation as resolved.
unsigned moveOrCopyTo = 0;
unsigned numberToMoveOrCopy = 1;
std::string newArgTypeName = "";
bool isPointerArg = false;
std::string defaultInitValue = "";
};

typedef std::map<unsigned, CastInfo> ArgCastMap;
Expand Down
80 changes: 78 additions & 2 deletions src/HipifyAction.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -91,6 +91,7 @@ const std::string sCudnnGetPoolingNdDescriptor = "cudnnGetPoolingNdDescriptor";
const std::string sCudnnSetLRNDescriptor = "cudnnSetLRNDescriptor";
const std::string sCudnnGetRNNDescriptor_v6 = "cudnnGetRNNDescriptor_v6";
const std::string sCudnnSetRNNDescriptor_v6 = "cudnnSetRNNDescriptor_v6";
const std::string sCudnnSetDropoutDescriptor = "cudnnSetDropoutDescriptor";
const std::string sCudnnSoftmaxForward = "cudnnSoftmaxForward";
const std::string sCudnnSoftmaxBackward = "cudnnSoftmaxBackward";
const std::string sCudnnConvolutionForward = "cudnnConvolutionForward";
Expand Down Expand Up @@ -263,6 +264,7 @@ std::string getCastType(hipify::CastTypes c) {
case e_add_var_argument: return "";
case e_move_argument: return "";
case e_replace_argument_with_const: return "";
case e_insert_new_argument: return "";
default: return "";
}
}
Expand Down Expand Up @@ -596,7 +598,8 @@ std::map<std::string, std::vector<ArgCastStruct>> FuncArgCasts {
{
{
{
{0, {e_remove_argument, cw_None}}
{0, {e_remove_argument, cw_None}},
{8, {e_insert_new_argument, cw_NeedsNewArgDecl, "hipify_biasMode", 0, 1, "miopenRNNBiasMode_t", true}}
},
true,
true
Expand All @@ -607,7 +610,21 @@ std::map<std::string, std::vector<ArgCastStruct>> FuncArgCasts {
{
{
{
{0, {e_remove_argument, cw_None}}
{0, {e_remove_argument, cw_None}},
{8, {e_insert_new_argument, cw_NeedsNewArgDecl, "hipify_biasMode", 0, 1, "miopenRNNBiasMode_t", false}}
},
true,
true
}
}
},
{sCudnnSetDropoutDescriptor,
{
{
{
{6, {e_insert_new_argument, cw_NeedsNewArgDecl, "hipify_use_mask", 6, 1, "bool", false}},
{7, {e_insert_new_argument, cw_NeedsNewArgDecl, "hipify_state_evo", 7, 1, "bool", false}},
{8, {e_insert_new_argument, cw_NeedsNewArgDecl, "hipify_rng_mode", 8, 1, "miopenRNGType_t", false}}
},
true,
true
Expand Down Expand Up @@ -2712,7 +2729,65 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result)
if (TranslateToRoc == true && cc.isToRoc == false) continue;
if (TranslateToRoc == false && cc.isToRoc == true && TranslateToMIOpen == false) continue;
clang::LangOptions DefaultLangOptions;
std::string combinedDeclText;
std::string combinedArgText;
bool hasInsertions = false;
std::vector<hipify::CastInfo> insertionWarnings;
clang::SourceLocation callBeginLoc = call->getBeginLoc();
unsigned callCol = SM.getSpellingColumnNumber(callBeginLoc);
clang::SourceLocation stmtInsertLoc = callBeginLoc.getLocWithOffset(-static_cast<int>(callCol - 1));
Comment thread
ranapratap55 marked this conversation as resolved.
std::string indent;
const char *indentPtr = SM.getCharacterData(stmtInsertLoc);
while (*indentPtr == ' ' || *indentPtr == '\t')
indent += *indentPtr++;
for (auto c : cc.castMap) {
if (c.second.castType == e_insert_new_argument) {
hasInsertions = true;
std::string initExpr = c.second.defaultInitValue.empty() ? "{}" : c.second.defaultInitValue;
combinedDeclText += indent + c.second.newArgTypeName + " " + c.second.constValToAddOrReplace + " = " + initExpr + ";\n";
std::string argText;
if (c.second.isPointerArg)
argText = "&" + c.second.constValToAddOrReplace;
else
argText = c.second.constValToAddOrReplace;
if (c.first < call->getNumArgs()) {
clang::SourceLocation argLoc = call->getArg(c.first)->getBeginLoc();
ct::Replacement middleRep(SM, argLoc, 0, argText + ", ");
clang::FullSourceLoc middleFullSL(argLoc, SM);
insertReplacement(middleRep, middleFullSL);
} else {
if (!combinedArgText.empty())
combinedArgText += ", ";
combinedArgText += argText;
}
insertionWarnings.push_back(c.second);
}
}
if (hasInsertions) {
ct::Replacement declRep(SM, stmtInsertLoc, 0, combinedDeclText);
clang::FullSourceLoc declFullSL(stmtInsertLoc, SM);
insertReplacement(declRep, declFullSL);
if (!combinedArgText.empty()) {
clang::SourceLocation insertLoc = call->getEndLoc();
std::string fullArgText = ", " + combinedArgText;
ct::Replacement argRep(SM, insertLoc, 0, fullArgText);
clang::FullSourceLoc argFullSL(insertLoc, SM);
insertReplacement(argRep, argFullSL);
}

for (auto &info : insertionWarnings) {
clang::DiagnosticsEngine &DE = getCompilerInstance().getDiagnostics();
const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning,
"HIP API '%0' requires additional argument '%1' of type '%2'. "
"A variable declaration has been inserted before the call. "
"Please initialize it appropriately.");
clang::FullSourceLoc warnFullSL(call->getBeginLoc(), SM);
DE.Report(warnFullSL, ID) << sName << info.constValToAddOrReplace << info.newArgTypeName;
}
}
for (auto c : cc.castMap) {
if (c.second.castType == e_insert_new_argument)
continue;
size_t length = 0;
unsigned int argNum = c.first;
clang::SmallString<40> XStr;
Expand Down Expand Up @@ -2991,6 +3066,7 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
sCudnnSetLRNDescriptor,
sCudnnGetRNNDescriptor_v6,
sCudnnSetRNNDescriptor_v6,
sCudnnSetDropoutDescriptor,
sCudnnSoftmaxForward,
sCudnnSoftmaxBackward,
sCudnnConvolutionForward,
Expand Down
6 changes: 4 additions & 2 deletions tests/unit_tests/synthetic/libraries/cudnn2miopen.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -745,10 +745,12 @@ int main() {
// CHECK: status = miopenGetDropoutDescriptor(DropoutDescriptor, handle, &dropout, &states, &seed);
status = cudnnGetDropoutDescriptor(DropoutDescriptor, handle, &dropout, &states, &seed);

// TODO [#837]: Insert bool use_mask, bool state_evo, miopenRNGType_t rng_mode in the hipified miopenGetDropoutDescriptor: will need variable declaration
// CUDA: cudnnStatus_t CUDNNWINAPI cudnnSetDropoutDescriptor(cudnnDropoutDescriptor_t dropoutDesc, cudnnHandle_t handle, float dropout, void* states, size_t stateSizeInBytes, unsigned long long seed);
// MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenSetDropoutDescriptor(miopenDropoutDescriptor_t dropoutDesc, miopenHandle_t handle, float dropout, void* states, size_t stateSizeInBytes, unsigned long long seed, bool use_mask, bool state_evo, miopenRNGType_t rng_mode);
// CHECK: status = miopenSetDropoutDescriptor(DropoutDescriptor, handle, dropout, states, reserveSpaceNumBytes, seed);
// CHECK: bool hipify_use_mask = {};
// CHECK-NEXT: bool hipify_state_evo = {};
// CHECK-NEXT: miopenRNGType_t hipify_rng_mode = {};
// CHECK-NEXT: status = miopenSetDropoutDescriptor(DropoutDescriptor, handle, dropout, states, reserveSpaceNumBytes, seed, hipify_use_mask, hipify_state_evo, hipify_rng_mode);
status = cudnnSetDropoutDescriptor(DropoutDescriptor, handle, dropout, states, reserveSpaceNumBytes, seed);

// TODO [#837]: Insert bool use_mask, bool state_evo, miopenRNGType_t rng_mode in the hipified miopenRestoreDropoutDescriptor: will need variable declaration
Expand Down
8 changes: 4 additions & 4 deletions tests/unit_tests/synthetic/libraries/cudnn2miopen_before_9000.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -61,16 +61,16 @@ int main() {
size_t reserveSpaceNumBytes = 0;

#if CUDNN_MAJOR < 9
// TODO [#837]: Insert miopenRNNBiasMode_t biasMode in the hipified miopenSetRNNDescriptor_V2 after miopenRNNMode_t rnnMode: will need variable declaration
// CUDA: CUDNN_DEPRECATED cudnnStatus_t CUDNNWINAPI cudnnSetRNNDescriptor_v6(cudnnHandle_t handle, cudnnRNNDescriptor_t rnnDesc, const int hiddenSize, const int numLayers, cudnnDropoutDescriptor_t dropoutDesc, cudnnRNNInputMode_t inputMode, cudnnDirectionMode_t direction, cudnnRNNMode_t cellMode, cudnnRNNAlgo_t algo, cudnnDataType_t mathPrec);
// MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenSetRNNDescriptor_V2(miopenRNNDescriptor_t rnnDesc, const int hsize, const int nlayers, miopenDropoutDescriptor_t dropoutDesc, miopenRNNInputMode_t inMode, miopenRNNDirectionMode_t direction, miopenRNNMode_t rnnMode, miopenRNNBiasMode_t biasMode, miopenRNNAlgo_t algo, miopenDataType_t dataType);
// CHECK: status = miopenSetRNNDescriptor_V2(RNNDescriptor, hiddenSize, layer, DropoutDescriptor, RNNInputMode, DirectionMode, RNNMode, RNNAlgo, dataType);
// CHECK: miopenRNNBiasMode_t hipify_biasMode = {};
// CHECK-NEXT: status = miopenSetRNNDescriptor_V2(RNNDescriptor, hiddenSize, layer, DropoutDescriptor, RNNInputMode, DirectionMode, RNNMode, hipify_biasMode, RNNAlgo, dataType);
status = cudnnSetRNNDescriptor_v6(handle, RNNDescriptor, hiddenSize, layer, DropoutDescriptor, RNNInputMode, DirectionMode, RNNMode, RNNAlgo, dataType);

// TODO [#837]: Insert miopenRNNBiasMode_t* biasMode in the hipified miopenGetRNNDescriptor_V2 after miopenRNNMode_t* rnnMode: will need variable declaration
// CUDA: CUDNN_DEPRECATED cudnnStatus_t CUDNNWINAPI cudnnGetRNNDescriptor_v6(cudnnHandle_t handle, cudnnRNNDescriptor_t rnnDesc, int* hiddenSize, int* numLayers, cudnnDropoutDescriptor_t* dropoutDesc, cudnnRNNInputMode_t* inputMode, cudnnDirectionMode_t* direction, cudnnRNNMode_t* cellMode, cudnnRNNAlgo_t* algo, cudnnDataType_t* mathPrec);
// MIOPEN: MIOPEN_EXPORT miopenStatus_t miopenGetRNNDescriptor_V2(miopenRNNDescriptor_t rnnDesc, int* hiddenSize, int* layer, miopenDropoutDescriptor_t* dropoutDesc, miopenRNNInputMode_t* inputMode, miopenRNNDirectionMode_t* dirMode, miopenRNNMode_t* rnnMode, miopenRNNBiasMode_t* biasMode, miopenRNNAlgo_t* algoMode, miopenDataType_t* dataType);
// CHECK: status = miopenGetRNNDescriptor_V2(RNNDescriptor, &hiddenSize, &layer, &DropoutDescriptor, &RNNInputMode, &DirectionMode, &RNNMode, &RNNAlgo, &dataType);
// CHECK: miopenRNNBiasMode_t hipify_biasMode = {};
// CHECK-NEXT: status = miopenGetRNNDescriptor_V2(RNNDescriptor, &hiddenSize, &layer, &DropoutDescriptor, &RNNInputMode, &DirectionMode, &RNNMode, &hipify_biasMode, &RNNAlgo, &dataType);
status = cudnnGetRNNDescriptor_v6(handle, RNNDescriptor, &hiddenSize, &layer, &DropoutDescriptor, &RNNInputMode, &DirectionMode, &RNNMode, &RNNAlgo, &dataType);

// CUDA: CUDNN_DEPRECATED cudnnStatus_t CUDNNWINAPI cudnnRNNBackwardWeights(cudnnHandle_t handle, const cudnnRNNDescriptor_t rnnDesc, const int seqLength, const cudnnTensorDescriptor_t* xDesc, const void* x, const cudnnTensorDescriptor_t hxDesc, const void* hx, const cudnnTensorDescriptor_t* yDesc, const void* y, const void* workSpace, size_t workSpaceSizeInBytes, const cudnnFilterDescriptor_t dwDesc, void* dw, const void* reserveSpace, size_t reserveSpaceSizeInBytes);
Expand Down