[HIPIFY][#837][#2062][feature][RFC][Partial] Implement e_insert_new_argument for CUDA-to-HIP APIs with variable declarations - Part 1

[ranapratap55]

Motivation

Issue #837, additional arguments must be declared before the function call. This is necessary because several cuDNN to MIOpen mappings require parameters such as use_mask, state_evo, and rng_mode for miopenSetDropoutDescriptor which lack CUDA equivalents and cannot be inferred from the current context.

Problem Statement

When hipify-clang encounters a CUDA API whose HIP/MIOpen equivalent has additional arguments that are not present in the original call, it currently has no mechanism to

1. Declare new variables before the function call
2. Append new variables as arguments after hipified call
3. Warn user that new variables need initialization before usage

Solution

Introduces e_insert_new_argument:

1. Declares a new variable before the function call
2. Appends the variable as an argument
3. Emits a warning to user for each inserted variable about intialization

Summary of changes in this PR:

• e_insert_new_argument type for APIs where MIOpen requires additional arguments not present in CUDA [HIPIFY][feature] Add a new function call transformation type "additional non-const arg" #837
• Variable declarations generated with value-initialization (= {}) and hipify_ prefix
• Middle insertions at correct positions, end-appends batched to work for same-offset constraint
• Compiler warning emitted for each generated variable instructing user to initialize appropriately

ToDo work (follow-up PRs):

• Braceless control flow scope: When API call is inside a single-statement if/for/while without braces, the inserted declarations break control flow. Fix: query AST parents to detect CompoundStmt and wrap with { }.
• Multiple-declaration guard: If the same API is called twice in same scope, the hipify_ variables would be declared twice. Fix: append a number (e.g., hipify_use_mask_1).

[emankov]

Good try on an important feature from our ToDo list, but see my comments, please.

[emankov]

The variable declaration logic determines the line start and inserts the new variables directly above the function call.

If the original function call is inside a single-statement control flow block without braces, this textual insertion will break the AST logic.

For instance:

if (status == CUDNN_STATUS_SUCCESS)
    cudnnSetDropoutDescriptor(...);

will be hipified to smth. like the following:

if (status == CUDNN_STATUS_SUCCESS)
bool use_mask;
bool state_evo;
miopenRNGType_t rng_mode;
    miopenSetDropoutDescriptor(...); // This is now outside the if-statement

[Recommendation]

To make this robust, check if the parent AST node is a CompoundStmt. If it's not (e.g., it's an IfStmt or ForStmt), you should generate { and } around the inserted declarations and the function call to preserve the original control flow scope.

[ranapratap55]

Correct, this is a real issue here.

I can implement the AST check in this current PR. It is not only to e_insert_new_argument but it applies to future feature that inserts statements above a call.

it can be designed as a reusable utility with multiple control flow variants like (if, else, for, while, do-while).

Would you prefer I handle it here, or as a dedicated follow-up PR?

[emankov]

The middle argument insertion is a very elegant solution!

Using call->getArg(c.first)->getBeginLoc() to seamlessly interleave the new arguments into the existing Clang replacement stream is much cleaner than reconstructing the entire argument list. This is the approach we prefer to use when dealing with arguments in Clang. Adding the {} default initialization is also a great touch to squash uninitialized variable warnings right out of the gate.

However, we need to address the scoping implementation, as it didn't fully resolve the control-flow vulnerability.

[ranapratap55]

The middle argument insertion is a very elegant solution!

Using call->getArg(c.first)->getBeginLoc() to seamlessly interleave the new arguments into the existing Clang replacement stream is much cleaner than reconstructing the entire argument list. This is the approach we prefer to use when dealing with arguments in Clang. Adding the {} default initialization is also a great touch to squash uninitialized variable warnings right out of the gate.

However, we need to address the scoping implementation, as it didn't fully resolve the control-flow vulnerability.

Marking this PR as partial and will create a ToDo for the below

1. Braceless control flow (AST parent)
2. Multiple declaration guard