Skip to content

[HIPIFY][#837][#2062][#2415] Add multi-declaration gaurd check for e_insert_new_argument - Part 2#2443

Open
ranapratap55 wants to merge 4 commits into
ROCm:amd-developfrom
ranapratap55:users/ranapratap55/non-linear-api-multi-declration
Open

[HIPIFY][#837][#2062][#2415] Add multi-declaration gaurd check for e_insert_new_argument - Part 2#2443
ranapratap55 wants to merge 4 commits into
ROCm:amd-developfrom
ranapratap55:users/ranapratap55/non-linear-api-multi-declration

Conversation

@ranapratap55
Copy link
Copy Markdown
Collaborator

Context

Follow-up PR to #2415 for e_insert_new_argument cast type that generates variable declarations above a function call when hipifying APIs where MIOpen requires additional arguments not present in CUDA.

One known limitation in #2415 was:

Multiple-declaration guard: If the same API is called twice in the same scope,
the hipify_ variables would be declared twice, causing a compile-time redefinition
error.

With this PR, we are fixing it.

Problem

Given a CUDA code with the same API called multiple times in the same scope:

status = cudnnSetDropoutDescriptor(desc, handle, dropout, states, size, seed);
status = cudnnSetDropoutDescriptor(desc, handle, dropout, states, size, seed);

#2415 would generate

bool hipify_use_mask = {};       // works
bool hipify_use_mask = {};       // fails — REDEFINITION ERROR

Solution

Added InsertedVarCounter that tracks the count for each base variable name that has been generated.

Also fixed warning emission loop with insertionVarNames vector so the compiler warning message correctly reports the unique name hipify_use_mask_1 instead of the base name hipify_use_mask.

With this change code would generate like:

bool hipify_use_mask = {};         // works
bool hipify_use_mask_1 = {};       // works - unique name
bool hipify_use_mask_2 = {};       // works - unique name

@ranapratap55 ranapratap55 self-assigned this Mar 24, 2026
@ranapratap55 ranapratap55 added enhancement Enhancement feature Feature request or implementation clang clang compiler related issue or change labels Mar 24, 2026
Comment thread src/HipifyAction.cpp Outdated
Comment thread src/HipifyAction.h Outdated
Comment thread src/HipifyAction.cpp Outdated
@ranapratap55 ranapratap55 force-pushed the users/ranapratap55/non-linear-api-multi-declration branch from 0136c06 to f29c773 Compare April 13, 2026 06:07
@ranapratap55 ranapratap55 requested a review from emankov April 13, 2026 06:10
Comment thread src/HipifyAction.h Outdated
// has been used per C++ block scope, to avoid redefinition when the same
// API is called multiple times in the same scope. getRawEncoding() is
// globally unique per block across the entire translation unit.
std::map<std::pair<unsigned, std::string>, unsigned> InsertedVarCounter;
Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

FileID is still needed to prevent loss of translation-unit isolation.
What about using std::tuple inside the map?

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done. Changed the map type to std::map<std::tuple<clang::FileID, unsigned, std::string>, unsigned>.

// CHECK-NEXT: status = miopenSetDropoutDescriptor(DropoutDescriptor, handle, dropout, states, reserveSpaceNumBytes, seed, hipify_use_mask_2, hipify_state_evo_2, hipify_rng_mode_2);
status = cudnnSetDropoutDescriptor(DropoutDescriptor, handle, dropout, states, reserveSpaceNumBytes, seed);

// CHECK: return 0;
Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Excessive check?

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Fixed.

Comment thread src/HipifyAction.cpp
Comment on lines 2778 to 2781
if (hasInsertions) {
ct::Replacement declRep(SM, stmtInsertLoc, 0, combinedDeclText);
clang::FullSourceLoc declFullSL(stmtInsertLoc, SM);
insertReplacement(declRep, declFullSL);
Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

[Assumption] The if-else construct without curly braces in the source CUDA code will lead to syntactically incorrect HIP code, for instance:

if (condition)
        status = cudnnSetDropoutDescriptor(desc, handle, dropout, states, size, seed);
    else
        status = CUDNN_STATUS_SUCCESS;

[ToDo] The if-else construct without curly braces should be presented in tessts.

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

yes, thanks. This is a ToDo 1 from #2415, Added a // TODO: block in .cpp and test file.

Comment thread src/HipifyAction.cpp Outdated
Comment on lines +2749 to +2751
if (const auto *cs = parents[0].get<clang::CompoundStmt>()) {
scopeKey = cs->getLBracLoc().getRawEncoding();
break;
Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

AST traverse here should cover all parents, not only the first one, which are very common in template Instantiations, macros, and implicit conversions.

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for this. Fixed in the latest revision.
Corrected to parent.get<clang::CompoundStmt>(), so every node at the current AST level is checked before ascending.

Comment thread src/HipifyAction.cpp Outdated
Comment on lines +2743 to +2749
for (auto c : cc.castMap) {
if (c.second.castType == e_insert_new_argument) {
const std::string &baseName = c.second.constValToAddOrReplace;
unsigned scopeKey = 0;
auto parents = Result.Context->getParents(*call);
while (!parents.empty()) {
if (const auto *cs = parents[0].get<clang::CompoundStmt>()) {
Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

[Performance] The expensive AST tree climb to find the CompoundStmt three separate times for the exact same function call.

Copy link
Copy Markdown
Collaborator Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks and fixed.
Moved the logic outside the for (auto c : cc.castMap) loop and AST climb happens exactly once.

@ranapratap55 ranapratap55 force-pushed the users/ranapratap55/non-linear-api-multi-declration branch from f29c773 to e75f434 Compare May 5, 2026 05:31
@ranapratap55 ranapratap55 requested a review from emankov May 7, 2026 05:28
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

clang clang compiler related issue or change enhancement Enhancement feature Feature request or implementation

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants