[HIPIFY][SWDEV-501193][RT][6.4.0] Added experimental support for `hip…

…GraphAddBatchMemOpNode` API + Updated synthetic tests, the regenerated `hipify-perl`, and `Driver` `CUDA2HIP` docs accordingly
emankov · Dec 11, 2024 · 05b7122 · 05b7122
1 parent 00ca4a0
commit 05b7122
Show file tree

Hide file tree

Showing 6 changed files with 59 additions and 17 deletions.
diff --git a/bin/hipify-perl b/bin/hipify-perl
@@ -1400,9 +1400,16 @@ my %removed_funcs = (
 my %experimental_funcs = (
     "cuStreamBatchMemOp_v2" => "6.4.0",
     "cuStreamBatchMemOp" => "6.4.0",
+    "cuGraphAddBatchMemOpNode" => "6.4.0",
     "CUstreamBatchMemOpParams_v1" => "6.4.0",
     "CUstreamBatchMemOpParams_union" => "6.4.0",
-    "CUstreamBatchMemOpParams" => "6.4.0"
+    "CUstreamBatchMemOpParams" => "6.4.0",
+    "CUDA_BATCH_MEM_OP_NODE_PARAMS_v2_st" => "6.4.0",
+    "CUDA_BATCH_MEM_OP_NODE_PARAMS_v2" => "6.4.0",
+    "CUDA_BATCH_MEM_OP_NODE_PARAMS_v1_st" => "6.4.0",
+    "CUDA_BATCH_MEM_OP_NODE_PARAMS_v1" => "6.4.0",
+    "CUDA_BATCH_MEM_OP_NODE_PARAMS_st" => "6.4.0",
+    "CUDA_BATCH_MEM_OP_NODE_PARAMS" => "6.4.0"
 );
 
 $print_stats = 1 if $examine;
@@ -1542,6 +1549,13 @@ sub subst {
 sub experimentalSubstitutions {
     subst("cuStreamBatchMemOp", "hipStreamBatchMemOp", "stream_memory");
     subst("cuStreamBatchMemOp_v2", "hipStreamBatchMemOp", "stream_memory");
+    subst("cuGraphAddBatchMemOpNode", "hipGraphAddBatchMemOpNode", "graph");
+    subst("CUDA_BATCH_MEM_OP_NODE_PARAMS", "hipBatchMemOpNodeParams", "type");
+    subst("CUDA_BATCH_MEM_OP_NODE_PARAMS_st", "hipBatchMemOpNodeParams", "type");
+    subst("CUDA_BATCH_MEM_OP_NODE_PARAMS_v1", "hipBatchMemOpNodeParams", "type");
+    subst("CUDA_BATCH_MEM_OP_NODE_PARAMS_v1_st", "hipBatchMemOpNodeParams", "type");
+    subst("CUDA_BATCH_MEM_OP_NODE_PARAMS_v2", "hipBatchMemOpNodeParams", "type");
+    subst("CUDA_BATCH_MEM_OP_NODE_PARAMS_v2_st", "hipBatchMemOpNodeParams", "type");
     subst("CUstreamBatchMemOpParams", "hipStreamBatchMemOpParams", "type");
     subst("CUstreamBatchMemOpParams_union", "hipStreamBatchMemOpParams_union", "type");
     subst("CUstreamBatchMemOpParams_v1", "hipStreamBatchMemOpParams", "type");
@@ -4076,7 +4090,6 @@ sub simpleSubstitutions {
     subst("cuDeviceGetGraphMemAttribute", "hipDeviceGetGraphMemAttribute", "graph");
     subst("cuDeviceGraphMemTrim", "hipDeviceGraphMemTrim", "graph");
     subst("cuDeviceSetGraphMemAttribute", "hipDeviceSetGraphMemAttribute", "graph");
-    subst("cuGraphAddBatchMemOpNode", "hipGraphAddBatchMemOpNode", "graph");
     subst("cuGraphAddChildGraphNode", "hipGraphAddChildGraphNode", "graph");
     subst("cuGraphAddDependencies", "hipGraphAddDependencies", "graph");
     subst("cuGraphAddEmptyNode", "hipGraphAddEmptyNode", "graph");

diff --git a/docs/tables/CUDA_Driver_API_functions_supported_by_HIP.md b/docs/tables/CUDA_Driver_API_functions_supported_by_HIP.md
@@ -30,12 +30,12 @@
 |`CUDA_ARRAY_SPARSE_PROPERTIES`|11.1| | | | | | | | | |
 |`CUDA_ARRAY_SPARSE_PROPERTIES_st`|11.1| | | | | | | | | |
 |`CUDA_ARRAY_SPARSE_PROPERTIES_v1`|11.3| | | | | | | | | |
-|`CUDA_BATCH_MEM_OP_NODE_PARAMS`|11.7| | | | | | | | | |
-|`CUDA_BATCH_MEM_OP_NODE_PARAMS_st`|11.7| | |12.2| | | | | | |
-|`CUDA_BATCH_MEM_OP_NODE_PARAMS_v1`|12.2| | | | | | | | | |
-|`CUDA_BATCH_MEM_OP_NODE_PARAMS_v1_st`|12.2| | | | | | | | | |
-|`CUDA_BATCH_MEM_OP_NODE_PARAMS_v2`|12.2| | | | | | | | | |
-|`CUDA_BATCH_MEM_OP_NODE_PARAMS_v2_st`|12.2| | | | | | | | | |
+|`CUDA_BATCH_MEM_OP_NODE_PARAMS`|11.7| | | |`hipBatchMemOpNodeParams`|6.4.0| | | |6.4.0|
+|`CUDA_BATCH_MEM_OP_NODE_PARAMS_st`|11.7| | |12.2|`hipBatchMemOpNodeParams`|6.4.0| | | |6.4.0|
+|`CUDA_BATCH_MEM_OP_NODE_PARAMS_v1`|12.2| | | |`hipBatchMemOpNodeParams`|6.4.0| | | |6.4.0|
+|`CUDA_BATCH_MEM_OP_NODE_PARAMS_v1_st`|12.2| | | |`hipBatchMemOpNodeParams`|6.4.0| | | |6.4.0|
+|`CUDA_BATCH_MEM_OP_NODE_PARAMS_v2`|12.2| | | |`hipBatchMemOpNodeParams`|6.4.0| | | |6.4.0|
+|`CUDA_BATCH_MEM_OP_NODE_PARAMS_v2_st`|12.2| | | |`hipBatchMemOpNodeParams`|6.4.0| | | |6.4.0|
 |`CUDA_CB`| | | | | | | | | | |
 |`CUDA_CHILD_GRAPH_NODE_PARAMS`|12.2| | | |`hipChildGraphNodeParams`|6.1.0| | | | |
 |`CUDA_CHILD_GRAPH_NODE_PARAMS_st`|12.2| | | |`hipChildGraphNodeParams`|6.1.0| | | | |
@@ -1889,7 +1889,7 @@
 |`cuDeviceGetGraphMemAttribute`|11.4| | | |`hipDeviceGetGraphMemAttribute`|5.3.0| | | | |
 |`cuDeviceGraphMemTrim`|11.4| | | |`hipDeviceGraphMemTrim`|5.3.0| | | | |
 |`cuDeviceSetGraphMemAttribute`|11.4| | | |`hipDeviceSetGraphMemAttribute`|5.3.0| | | | |
-|`cuGraphAddBatchMemOpNode`|11.7| | | |`hipGraphAddBatchMemOpNode`| | | | | |
+|`cuGraphAddBatchMemOpNode`|11.7| | | |`hipGraphAddBatchMemOpNode`|6.4.0| | | |6.4.0|
 |`cuGraphAddChildGraphNode`|10.0| | | |`hipGraphAddChildGraphNode`|5.0.0| | | | |
 |`cuGraphAddDependencies`|10.0| | | |`hipGraphAddDependencies`|4.5.0| | | | |
 |`cuGraphAddDependencies_v2`|12.3| | | | | | | | | |

diff --git a/src/CUDA2HIP_Driver_API_functions.cpp b/src/CUDA2HIP_Driver_API_functions.cpp
@@ -808,8 +808,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_FUNCTION_MAP {
   {"cuDeviceSetGraphMemAttribute",                                {"hipDeviceSetGraphMemAttribute",                               "", CONV_GRAPH, API_DRIVER, SEC::GRAPH}},
   // cudaGraphInstantiateWithFlags
   {"cuGraphInstantiateWithFlags",                                 {"hipGraphInstantiateWithFlags",                                "", CONV_GRAPH, API_DRIVER, SEC::GRAPH}},
-  //
-  {"cuGraphAddBatchMemOpNode",                                    {"hipGraphAddBatchMemOpNode",                                   "", CONV_GRAPH, API_DRIVER, SEC::GRAPH}},
+  // no analogue yet
+  {"cuGraphAddBatchMemOpNode",                                    {"hipGraphAddBatchMemOpNode",                                   "", CONV_GRAPH, API_DRIVER, SEC::GRAPH, HIP_EXPERIMENTAL}},
   //
   {"cuGraphBatchMemOpNodeGetParams",                              {"hipGraphBatchMemOpNodeGetParams",                             "", CONV_GRAPH, API_DRIVER, SEC::GRAPH}},
   //
@@ -1666,6 +1666,7 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_DRIVER_FUNCTION_VER_MAP {
   {"hipDrvGraphExecMemcpyNodeSetParams",                          {HIP_6030, HIP_0,    HIP_0,  }},
   {"hipDrvGraphExecMemsetNodeSetParams",                          {HIP_6030, HIP_0,    HIP_0,  }},
   {"hipStreamBatchMemOp",                                         {HIP_6040, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"hipGraphAddBatchMemOpNode",                                   {HIP_6040, HIP_0,    HIP_0,  HIP_LATEST}},
 };
 
 const std::map<llvm::StringRef, cudaAPIChangedVersions> CUDA_DRIVER_FUNCTION_CHANGED_VER_MAP {

diff --git a/src/CUDA2HIP_Driver_API_types.cpp b/src/CUDA2HIP_Driver_API_types.cpp
@@ -357,12 +357,12 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP {
   {"memoryBarrier",                                                    {"hipMemoryBarrier",                                         "", CONV_TYPE, API_DRIVER, SEC::DATA_TYPES, HIP_UNSUPPORTED}},
 
   // no analogue
-  {"CUDA_BATCH_MEM_OP_NODE_PARAMS_st",                                 {"HIP_BATCH_MEM_OP_NODE_PARAMS",                             "", CONV_TYPE, API_DRIVER, SEC::DATA_TYPES, HIP_UNSUPPORTED | CUDA_REMOVED}},
-  {"CUDA_BATCH_MEM_OP_NODE_PARAMS_v1_st",                              {"HIP_BATCH_MEM_OP_NODE_PARAMS",                             "", CONV_TYPE, API_DRIVER, SEC::DATA_TYPES, HIP_UNSUPPORTED}},
-  {"CUDA_BATCH_MEM_OP_NODE_PARAMS_v2_st",                              {"HIP_BATCH_MEM_OP_NODE_PARAMS_v2",                          "", CONV_TYPE, API_DRIVER, SEC::DATA_TYPES, HIP_UNSUPPORTED}},
-  {"CUDA_BATCH_MEM_OP_NODE_PARAMS",                                    {"HIP_BATCH_MEM_OP_NODE_PARAMS",                             "", CONV_TYPE, API_DRIVER, SEC::DATA_TYPES, HIP_UNSUPPORTED}},
-  {"CUDA_BATCH_MEM_OP_NODE_PARAMS_v1",                                 {"HIP_BATCH_MEM_OP_NODE_PARAMS",                             "", CONV_TYPE, API_DRIVER, SEC::DATA_TYPES, HIP_UNSUPPORTED}},
-  {"CUDA_BATCH_MEM_OP_NODE_PARAMS_v2",                                 {"HIP_BATCH_MEM_OP_NODE_PARAMS_v2",                          "", CONV_TYPE, API_DRIVER, SEC::DATA_TYPES, HIP_UNSUPPORTED}},
+  {"CUDA_BATCH_MEM_OP_NODE_PARAMS_st",                                 {"hipBatchMemOpNodeParams",                                  "", CONV_TYPE, API_DRIVER, SEC::DATA_TYPES, HIP_EXPERIMENTAL | CUDA_REMOVED}},
+  {"CUDA_BATCH_MEM_OP_NODE_PARAMS_v1_st",                              {"hipBatchMemOpNodeParams",                                  "", CONV_TYPE, API_DRIVER, SEC::DATA_TYPES, HIP_EXPERIMENTAL}},
+  {"CUDA_BATCH_MEM_OP_NODE_PARAMS_v2_st",                              {"hipBatchMemOpNodeParams",                                  "", CONV_TYPE, API_DRIVER, SEC::DATA_TYPES, HIP_EXPERIMENTAL}},
+  {"CUDA_BATCH_MEM_OP_NODE_PARAMS",                                    {"hipBatchMemOpNodeParams",                                  "", CONV_TYPE, API_DRIVER, SEC::DATA_TYPES, HIP_EXPERIMENTAL}},
+  {"CUDA_BATCH_MEM_OP_NODE_PARAMS_v1",                                 {"hipBatchMemOpNodeParams",                                  "", CONV_TYPE, API_DRIVER, SEC::DATA_TYPES, HIP_EXPERIMENTAL}},
+  {"CUDA_BATCH_MEM_OP_NODE_PARAMS_v2",                                 {"hipBatchMemOpNodeParams",                                  "", CONV_TYPE, API_DRIVER, SEC::DATA_TYPES, HIP_EXPERIMENTAL}},
 
   // cudaLaunchAttribute_st
   {"CUlaunchAttribute_st",                                             {"hipLaunchAttribute",                                       "", CONV_TYPE, API_DRIVER, SEC::DATA_TYPES, HIP_UNSUPPORTED}},
@@ -4306,4 +4306,5 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_DRIVER_TYPE_NAME_VER_MAP {
   {"hipStreamLegacy",                                                  {HIP_6020, HIP_0,    HIP_0,  }},
   {"hipStreamBatchMemOpParams_union",                                  {HIP_6040, HIP_0,    HIP_0,  HIP_LATEST}},
   {"hipStreamBatchMemOpParams",                                        {HIP_6040, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"hipBatchMemOpNodeParams",                                          {HIP_6040, HIP_0,    HIP_0,  HIP_LATEST}},
 };
diff --git a/tests/unit_tests/synthetic/driver_functions.cu b/tests/unit_tests/synthetic/driver_functions.cu
@@ -1877,6 +1877,14 @@ int main() {
   // HIP: hipError_t hipStreamBatchMemOp(hipStream_t stream, unsigned int count, hipStreamBatchMemOpParams* paramArray, unsigned int flags);
   // CHECK: result = hipStreamBatchMemOp(stream, icount, &streamBatchMemOpParams, flags);
   result = cuStreamBatchMemOp_v2(stream, icount, &streamBatchMemOpParams, flags);
+
+  // CHECK: hipBatchMemOpNodeParams BATCH_MEM_OP_NODE_PARAMS;
+  CUDA_BATCH_MEM_OP_NODE_PARAMS BATCH_MEM_OP_NODE_PARAMS;
+
+  // CUDA: CUresult CUDAAPI cuGraphAddBatchMemOpNode(CUgraphNode *phGraphNode, CUgraph hGraph, const CUgraphNode *dependencies, size_t numDependencies, const CUDA_BATCH_MEM_OP_NODE_PARAMS *nodeParams);
+  // HIP: hipError_t hipGraphAddBatchMemOpNode(hipGraphNode_t *phGraphNode, hipGraph_t hGraph, const hipGraphNode_t* dependencies, size_t numDependencies, const hipBatchMemOpNodeParams* nodeParams);
+  // CHECK: result = hipGraphAddBatchMemOpNode(&graphNode, graph, &graphNode2, bytes, &BATCH_MEM_OP_NODE_PARAMS);
+  result = cuGraphAddBatchMemOpNode(&graphNode, graph, &graphNode2, bytes, &BATCH_MEM_OP_NODE_PARAMS);
 #endif
 
 #if CUDA_VERSION >= 12000

diff --git a/tests/unit_tests/synthetic/driver_structs.cu b/tests/unit_tests/synthetic/driver_structs.cu
@@ -308,6 +308,16 @@ int main() {
   CUDA_MEM_ALLOC_NODE_PARAMS_st MEM_ALLOC_NODE_PARAMS_st;
 #endif
 
+#if CUDA_VERSION >= 11070
+  // CHECK: hipBatchMemOpNodeParams BATCH_MEM_OP_NODE_PARAMS;
+  CUDA_BATCH_MEM_OP_NODE_PARAMS BATCH_MEM_OP_NODE_PARAMS;
+#endif
+
+#if CUDA_VERSION >= 11070 && CUDA_VERSION < 12020
+  // CHECK: hipBatchMemOpNodeParams BATCH_MEM_OP_NODE_PARAMS_st;
+  CUDA_BATCH_MEM_OP_NODE_PARAMS_st BATCH_MEM_OP_NODE_PARAMS_st;
+#endif
+
 #if CUDA_VERSION >= 12000
   // CHECK: hipGraphInstantiateParams GRAPH_INSTANTIATE_PARAMS_st;
   // CHECK-NEXT: hipGraphInstantiateParams GRAPH_INSTANTIATE_PARAMS;
@@ -360,6 +370,15 @@ int main() {
   // CHECK-NEXT: hipGraphNodeParams graphNodeParams;
   CUgraphNodeParams_st graphNodeParams_st;
   CUgraphNodeParams graphNodeParams;
+
+  // CHECK: hipBatchMemOpNodeParams BATCH_MEM_OP_NODE_PARAMS_v1_st;
+  // CHECK-NEXT: hipBatchMemOpNodeParams BATCH_MEM_OP_NODE_PARAMS_v1;
+  // CHECK-NEXT: hipBatchMemOpNodeParams BATCH_MEM_OP_NODE_PARAMS_v2_st;
+  // CHECK-NEXT: hipBatchMemOpNodeParams BATCH_MEM_OP_NODE_PARAMS_v2;
+  CUDA_BATCH_MEM_OP_NODE_PARAMS_v1_st BATCH_MEM_OP_NODE_PARAMS_v1_st;
+  CUDA_BATCH_MEM_OP_NODE_PARAMS_v1 BATCH_MEM_OP_NODE_PARAMS_v1;
+  CUDA_BATCH_MEM_OP_NODE_PARAMS_v2_st BATCH_MEM_OP_NODE_PARAMS_v2_st;
+  CUDA_BATCH_MEM_OP_NODE_PARAMS_v2 BATCH_MEM_OP_NODE_PARAMS_v2;
 #endif
 
 #if CUDA_VERSION >= 12030