diff --git a/bin/hipify-perl b/bin/hipify-perl index 1d65632a..bd422769 100755 --- 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 index 4107fb85..a2559132 100644 --- 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 index b482e20e..6b9850d4 100644 --- a/src/CUDA2HIP_Driver_API_functions.cpp +++ b/src/CUDA2HIP_Driver_API_functions.cpp @@ -808,8 +808,8 @@ const std::map 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 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 CUDA_DRIVER_FUNCTION_CHANGED_VER_MAP { diff --git a/src/CUDA2HIP_Driver_API_types.cpp b/src/CUDA2HIP_Driver_API_types.cpp index 7277c671..1d77e091 100644 --- a/src/CUDA2HIP_Driver_API_types.cpp +++ b/src/CUDA2HIP_Driver_API_types.cpp @@ -357,12 +357,12 @@ const std::map 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 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 index 4fb6ea60..810c1d44 100644 --- 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 index 8f9201e9..476c8d9d 100644 --- 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