[HIPIFY] cudaMemcpy2DFromArray(Async) support

[ROCm/clr commit: 5dd00bdf52]

projects/clr/hipamd/bin/hipify-perl

@@ -269,6 +269,8 @@ sub simpleSubstitutions {
     $ft{'memory'} += s/\bcudaMemcpy\b/hipMemcpy/g;
     $ft{'memory'} += s/\bcudaMemcpy2D\b/hipMemcpy2D/g;
     $ft{'memory'} += s/\bcudaMemcpy2DAsync\b/hipMemcpy2DAsync/g;
+    $ft{'memory'} += s/\bcudaMemcpy2DFromArray\b/hipMemcpy2DFromArray/g;
+    $ft{'memory'} += s/\bcudaMemcpy2DFromArrayAsync\b/hipMemcpy2DFromArrayAsync/g;
     $ft{'memory'} += s/\bcudaMemcpy2DToArray\b/hipMemcpy2DToArray/g;
     $ft{'memory'} += s/\bcudaMemcpy3D\b/hipMemcpy3D/g;
     $ft{'memory'} += s/\bcudaMemcpy3DAsync\b/hipMemcpy3DAsync/g;

projects/clr/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md

@@ -160,8 +160,8 @@
 | `cudaMemcpy2D`                                            | `hipMemcpy2D`                 |
 | `cudaMemcpy2DArrayToArray`                                |                               |
 | `cudaMemcpy2DAsync`                                       | `hipMemcpy2DAsync`            |
-| `cudaMemcpy2DFromArray`                                   |                               |
-| `cudaMemcpy2DFromArrayAsync`                              |                               |
+| `cudaMemcpy2DFromArray`                                   | `hipMemcpy2DFromArray`        |
+| `cudaMemcpy2DFromArrayAsync`                              | `hipMemcpy2DFromArrayAsync`   |
 | `cudaMemcpy2DToArray`                                     | `hipMemcpy2DToArray`          |
 | `cudaMemcpy2DToArrayAsync`                                |                               |
 | `cudaMemcpy3D`                                            | `hipMemcpy3D`                 |

projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp

@@ -293,9 +293,9 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_FUNCTION_MAP{
   // NOTE: Not equal to cuMemcpy2DAsync due to different signatures
   {"cudaMemcpy2DAsync",                                       {"hipMemcpy2DAsync",                                       "", CONV_MEMORY, API_RUNTIME}},
   // no analogue
-  {"cudaMemcpy2DFromArray",                                   {"hipMemcpy2DFromArray",                                   "", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
+  {"cudaMemcpy2DFromArray",                                   {"hipMemcpy2DFromArray",                                   "", CONV_MEMORY, API_RUNTIME}},
   // no analogue
-  {"cudaMemcpy2DFromArrayAsync",                              {"hipMemcpy2DFromArrayAsync",                              "", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}},
+  {"cudaMemcpy2DFromArrayAsync",                              {"hipMemcpy2DFromArrayAsync",                              "", CONV_MEMORY, API_RUNTIME}},
   // no analogue
   {"cudaMemcpy2DToArray",                                     {"hipMemcpy2DToArray",                                     "", CONV_MEMORY, API_RUNTIME}},
   // no analogue