[HIPIFY][#837][#2062][#2415] Add multi-declaration gaurd check for e_insert_new_argument - Part 2#2443
Conversation
0136c06 to
f29c773
Compare
| // 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; |
There was a problem hiding this comment.
FileID is still needed to prevent loss of translation-unit isolation.
What about using std::tuple inside the map?
There was a problem hiding this comment.
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; |
| if (hasInsertions) { | ||
| ct::Replacement declRep(SM, stmtInsertLoc, 0, combinedDeclText); | ||
| clang::FullSourceLoc declFullSL(stmtInsertLoc, SM); | ||
| insertReplacement(declRep, declFullSL); |
There was a problem hiding this comment.
[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.
There was a problem hiding this comment.
yes, thanks. This is a ToDo 1 from #2415, Added a // TODO: block in .cpp and test file.
| if (const auto *cs = parents[0].get<clang::CompoundStmt>()) { | ||
| scopeKey = cs->getLBracLoc().getRawEncoding(); | ||
| break; |
There was a problem hiding this comment.
AST traverse here should cover all parents, not only the first one, which are very common in template Instantiations, macros, and implicit conversions.
There was a problem hiding this comment.
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.
| 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>()) { |
There was a problem hiding this comment.
[Performance] The expensive AST tree climb to find the CompoundStmt three separate times for the exact same function call.
There was a problem hiding this comment.
Thanks and fixed.
Moved the logic outside the for (auto c : cc.castMap) loop and AST climb happens exactly once.
…VarCounter per file counter
… multi-function tests
f29c773 to
e75f434
Compare
Context
Follow-up PR to #2415 for
e_insert_new_argumentcast 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:
With this PR, we are fixing it.
Problem
Given a CUDA code with the same API called multiple times in the same scope:
#2415 would generate
Solution
Added
InsertedVarCounterthat tracks the count for each base variable name that has been generated.Also fixed warning emission loop with
insertionVarNamesvector so the compiler warning message correctly reports the unique namehipify_use_mask_1instead of the base namehipify_use_mask.With this change code would generate like: