[HIPIFY][rocRAND][feature] Support for cuRAND -> rocRAND hipificati…

…on - Step 10 - Functions + Updated synthetic tests for `rocRAND` and `hipRAND`, the regenerated `hipify-perl`, and `RAND` `CUDA2HIP` documentation
emankov · Oct 1, 2024 · b72111c · b72111c
1 parent 0782547
commit b72111c
Show file tree

Hide file tree

Showing 7 changed files with 149 additions and 40 deletions.
diff --git a/bin/hipify-perl b/bin/hipify-perl
@@ -2143,6 +2143,8 @@ sub rocSubstitutions {
     subst("cudnnTransformTensor", "miopenTransformTensor", "library");
     subst("curandCreateGenerator", "rocrand_create_generator", "library");
     subst("curandCreateGeneratorHost", "rocrand_create_generator_host_blocking", "library");
+    subst("curandCreatePoissonDistribution", "rocrand_create_poisson_distribution", "library");
+    subst("curandDestroyDistribution", "rocrand_destroy_discrete_distribution", "library");
     subst("curandDestroyGenerator", "rocrand_destroy_generator", "library");
     subst("curandGenerate", "rocrand_generate", "library");
     subst("curandGenerateLogNormal", "rocrand_generate_log_normal", "library");
@@ -2154,8 +2156,17 @@ sub rocSubstitutions {
     subst("curandGenerateSeeds", "rocrand_initialize_generator", "library");
     subst("curandGenerateUniform", "rocrand_generate_uniform", "library");
     subst("curandGenerateUniformDouble", "rocrand_generate_uniform_double", "library");
+    subst("curandGetDirectionVectors32", "rocrand_get_direction_vectors32", "library");
+    subst("curandGetDirectionVectors64", "rocrand_get_direction_vectors64", "library");
+    subst("curandGetScrambleConstants32", "rocrand_get_scramble_constants32", "library");
+    subst("curandGetScrambleConstants64", "rocrand_get_scramble_constants64", "library");
+    subst("curandGetVersion", "rocrand_get_version", "library");
+    subst("curandMakeMTGP32Constants", "rocrand_make_constant", "library");
+    subst("curandMakeMTGP32KernelState", "rocrand_make_state_mtgp32", "library");
     subst("curandSetGeneratorOffset", "rocrand_set_offset", "library");
+    subst("curandSetGeneratorOrdering", "rocrand_set_ordering", "library");
     subst("curandSetPseudoRandomGeneratorSeed", "rocrand_set_seed", "library");
+    subst("curandSetQuasiRandomGeneratorDimensions", "rocrand_set_quasi_random_generator_dimensions", "library");
     subst("curandSetStream", "rocrand_set_stream", "library");
     subst("cusolverDnCpotrf", "rocsolver_cpotrf", "library");
     subst("cusolverDnCreate", "rocblas_create_handle", "library");
@@ -2764,6 +2775,8 @@ sub rocSubstitutions {
     subst("cudnnTensorFormat_t", "miopenTensorLayout_t", "type");
     subst("curandDirectionVectorSet", "rocrand_direction_vector_set", "type");
     subst("curandDirectionVectorSet_t", "rocrand_direction_vector_set", "type");
+    subst("curandDiscreteDistribution_st", "rocrand_discrete_distribution_st", "type");
+    subst("curandDiscreteDistribution_t", "rocrand_discrete_distribution", "type");
     subst("curandGenerator_st", "rocrand_generator_base_type", "type");
     subst("curandGenerator_t", "rocrand_generator", "type");
     subst("curandOrdering", "rocrand_ordering", "type");

diff --git a/docs/tables/CURAND_API_supported_by_HIP_and_ROC.md b/docs/tables/CURAND_API_supported_by_HIP_and_ROC.md
@@ -57,8 +57,8 @@
 |`curandDirectionVectorSet_t`| | | | |`hiprandDirectionVectorSet_t`|6.0.0| | | | |`rocrand_direction_vector_set`|6.0.0| | | | |
 |`curandDirectionVectors32_t`| | | | |`hiprandDirectionVectors32_t`|1.5.0| | | | | | | | | | |
 |`curandDirectionVectors64_t`| | | | |`hiprandDirectionVectors64_t`|6.0.0| | | | | | | | | | |
-|`curandDiscreteDistribution_st`| | | | |`hiprandDiscreteDistribution_st`|1.5.0| | | | | | | | | | |
-|`curandDiscreteDistribution_t`| | | | |`hiprandDiscreteDistribution_t`|1.5.0| | | | | | | | | | |
+|`curandDiscreteDistribution_st`| | | | |`hiprandDiscreteDistribution_st`|1.5.0| | | | |`rocrand_discrete_distribution_st`|1.5.0| | | | |
+|`curandDiscreteDistribution_t`| | | | |`hiprandDiscreteDistribution_t`|1.5.0| | | | |`rocrand_discrete_distribution`|1.5.0| | | | |
 |`curandDistributionM2Shift_st`| | | | | | | | | | | | | | | | |
 |`curandDistributionM2Shift_t`| | | | | | | | | | | | | | | | |
 |`curandDistributionShift_st`| | | | | | | | | | | | | | | | |
@@ -106,8 +106,8 @@
 |:--|:-:|:-:|:-:|:-:|:--|:-:|:-:|:-:|:-:|:-:|:--|:-:|:-:|:-:|:-:|:-:|
 |`curandCreateGenerator`| | | | |`hiprandCreateGenerator`|1.5.0| | | | |`rocrand_create_generator`|1.5.0| | | | |
 |`curandCreateGeneratorHost`| | | | |`hiprandCreateGeneratorHost`|1.5.0| | | | |`rocrand_create_generator_host_blocking`|6.2.0| | | | |
-|`curandCreatePoissonDistribution`| | | | |`hiprandCreatePoissonDistribution`|1.5.0| | | | | | | | | | |
-|`curandDestroyDistribution`| | | | |`hiprandDestroyDistribution`|1.5.0| | | | | | | | | | |
+|`curandCreatePoissonDistribution`| | | | |`hiprandCreatePoissonDistribution`|1.5.0| | | | |`rocrand_create_poisson_distribution`|1.5.0| | | | |
+|`curandDestroyDistribution`| | | | |`hiprandDestroyDistribution`|1.5.0| | | | |`rocrand_destroy_discrete_distribution`|1.5.0| | | | |
 |`curandDestroyGenerator`| | | | |`hiprandDestroyGenerator`|1.5.0| | | | |`rocrand_destroy_generator`|1.5.0| | | | |
 |`curandGenerate`| | | | |`hiprandGenerate`|1.5.0| | | | |`rocrand_generate`|1.5.0| | | | |
 |`curandGenerateLogNormal`| | | | |`hiprandGenerateLogNormal`|1.5.0| | | | |`rocrand_generate_log_normal`|1.5.0| | | | |
@@ -119,18 +119,18 @@
 |`curandGenerateSeeds`| | | | |`hiprandGenerateSeeds`|1.5.0| | | | |`rocrand_initialize_generator`|1.5.0| | | | |
 |`curandGenerateUniform`| | | | |`hiprandGenerateUniform`|1.5.0| | | | |`rocrand_generate_uniform`|1.5.0| | | | |
 |`curandGenerateUniformDouble`| | | | |`hiprandGenerateUniformDouble`|1.5.0| | | | |`rocrand_generate_uniform_double`|1.5.0| | | | |
-|`curandGetDirectionVectors32`| | | | |`hiprandGetDirectionVectors32`|6.0.0| | | | | | | | | | |
-|`curandGetDirectionVectors64`| | | | |`hiprandGetDirectionVectors64`|6.0.0| | | | | | | | | | |
+|`curandGetDirectionVectors32`| | | | |`hiprandGetDirectionVectors32`|6.0.0| | | | |`rocrand_get_direction_vectors32`|6.0.0| | | | |
+|`curandGetDirectionVectors64`| | | | |`hiprandGetDirectionVectors64`|6.0.0| | | | |`rocrand_get_direction_vectors64`|6.0.0| | | | |
 |`curandGetProperty`|8.0| | | | | | | | | | | | | | | |
-|`curandGetScrambleConstants32`| | | | |`hiprandGetScrambleConstants32`|6.0.0| | | | | | | | | | |
-|`curandGetScrambleConstants64`| | | | |`hiprandGetScrambleConstants64`|6.0.0| | | | | | | | | | |
-|`curandGetVersion`| | | | |`hiprandGetVersion`|1.5.0| | | | | | | | | | |
-|`curandMakeMTGP32Constants`| | | | |`hiprandMakeMTGP32Constants`|1.5.0| | | | | | | | | | |
-|`curandMakeMTGP32KernelState`| | | | |`hiprandMakeMTGP32KernelState`|1.5.0| | | | | | | | | | |
+|`curandGetScrambleConstants32`| | | | |`hiprandGetScrambleConstants32`|6.0.0| | | | |`rocrand_get_scramble_constants32`|6.0.0| | | | |
+|`curandGetScrambleConstants64`| | | | |`hiprandGetScrambleConstants64`|6.0.0| | | | |`rocrand_get_scramble_constants64`|6.0.0| | | | |
+|`curandGetVersion`| | | | |`hiprandGetVersion`|1.5.0| | | | |`rocrand_get_version`|1.5.0| | | | |
+|`curandMakeMTGP32Constants`| | | | |`hiprandMakeMTGP32Constants`|1.5.0| | | | |`rocrand_make_constant`|1.5.0| | | | |
+|`curandMakeMTGP32KernelState`| | | | |`hiprandMakeMTGP32KernelState`|1.5.0| | | | |`rocrand_make_state_mtgp32`|1.5.0| | | | |
 |`curandSetGeneratorOffset`| | | | |`hiprandSetGeneratorOffset`|1.5.0| | | | |`rocrand_set_offset`|1.5.0| | | | |
-|`curandSetGeneratorOrdering`| | | | |`hiprandSetGeneratorOrdering`|6.2.0| | | | | | | | | | |
+|`curandSetGeneratorOrdering`| | | | |`hiprandSetGeneratorOrdering`|6.2.0| | | | |`rocrand_set_ordering`|5.5.0| | | | |
 |`curandSetPseudoRandomGeneratorSeed`| | | | |`hiprandSetPseudoRandomGeneratorSeed`|1.5.0| | | | |`rocrand_set_seed`|1.5.0| | | | |
-|`curandSetQuasiRandomGeneratorDimensions`| | | | |`hiprandSetQuasiRandomGeneratorDimensions`|1.5.0| | | | | | | | | | |
+|`curandSetQuasiRandomGeneratorDimensions`| | | | |`hiprandSetQuasiRandomGeneratorDimensions`|1.5.0| | | | |`rocrand_set_quasi_random_generator_dimensions`|1.5.0| | | | |
 |`curandSetStream`| | | | |`hiprandSetStream`|1.5.0| | | | |`rocrand_set_stream`|1.5.0| | | | |
 
 ## **3. Device API Functions**

diff --git a/docs/tables/CURAND_API_supported_by_ROC.md b/docs/tables/CURAND_API_supported_by_ROC.md
@@ -57,8 +57,8 @@
 |`curandDirectionVectorSet_t`| | | | |`rocrand_direction_vector_set`|6.0.0| | | | |
 |`curandDirectionVectors32_t`| | | | | | | | | | |
 |`curandDirectionVectors64_t`| | | | | | | | | | |
-|`curandDiscreteDistribution_st`| | | | | | | | | | |
-|`curandDiscreteDistribution_t`| | | | | | | | | | |
+|`curandDiscreteDistribution_st`| | | | |`rocrand_discrete_distribution_st`|1.5.0| | | | |
+|`curandDiscreteDistribution_t`| | | | |`rocrand_discrete_distribution`|1.5.0| | | | |
 |`curandDistributionM2Shift_st`| | | | | | | | | | |
 |`curandDistributionM2Shift_t`| | | | | | | | | | |
 |`curandDistributionShift_st`| | | | | | | | | | |
@@ -106,8 +106,8 @@
 |:--|:-:|:-:|:-:|:-:|:--|:-:|:-:|:-:|:-:|:-:|
 |`curandCreateGenerator`| | | | |`rocrand_create_generator`|1.5.0| | | | |
 |`curandCreateGeneratorHost`| | | | |`rocrand_create_generator_host_blocking`|6.2.0| | | | |
-|`curandCreatePoissonDistribution`| | | | | | | | | | |
-|`curandDestroyDistribution`| | | | | | | | | | |
+|`curandCreatePoissonDistribution`| | | | |`rocrand_create_poisson_distribution`|1.5.0| | | | |
+|`curandDestroyDistribution`| | | | |`rocrand_destroy_discrete_distribution`|1.5.0| | | | |
 |`curandDestroyGenerator`| | | | |`rocrand_destroy_generator`|1.5.0| | | | |
 |`curandGenerate`| | | | |`rocrand_generate`|1.5.0| | | | |
 |`curandGenerateLogNormal`| | | | |`rocrand_generate_log_normal`|1.5.0| | | | |
@@ -119,18 +119,18 @@
 |`curandGenerateSeeds`| | | | |`rocrand_initialize_generator`|1.5.0| | | | |
 |`curandGenerateUniform`| | | | |`rocrand_generate_uniform`|1.5.0| | | | |
 |`curandGenerateUniformDouble`| | | | |`rocrand_generate_uniform_double`|1.5.0| | | | |
-|`curandGetDirectionVectors32`| | | | | | | | | | |
-|`curandGetDirectionVectors64`| | | | | | | | | | |
+|`curandGetDirectionVectors32`| | | | |`rocrand_get_direction_vectors32`|6.0.0| | | | |
+|`curandGetDirectionVectors64`| | | | |`rocrand_get_direction_vectors64`|6.0.0| | | | |
 |`curandGetProperty`|8.0| | | | | | | | | |
-|`curandGetScrambleConstants32`| | | | | | | | | | |
-|`curandGetScrambleConstants64`| | | | | | | | | | |
-|`curandGetVersion`| | | | | | | | | | |
-|`curandMakeMTGP32Constants`| | | | | | | | | | |
-|`curandMakeMTGP32KernelState`| | | | | | | | | | |
+|`curandGetScrambleConstants32`| | | | |`rocrand_get_scramble_constants32`|6.0.0| | | | |
+|`curandGetScrambleConstants64`| | | | |`rocrand_get_scramble_constants64`|6.0.0| | | | |
+|`curandGetVersion`| | | | |`rocrand_get_version`|1.5.0| | | | |
+|`curandMakeMTGP32Constants`| | | | |`rocrand_make_constant`|1.5.0| | | | |
+|`curandMakeMTGP32KernelState`| | | | |`rocrand_make_state_mtgp32`|1.5.0| | | | |
 |`curandSetGeneratorOffset`| | | | |`rocrand_set_offset`|1.5.0| | | | |
-|`curandSetGeneratorOrdering`| | | | | | | | | | |
+|`curandSetGeneratorOrdering`| | | | |`rocrand_set_ordering`|5.5.0| | | | |
 |`curandSetPseudoRandomGeneratorSeed`| | | | |`rocrand_set_seed`|1.5.0| | | | |
-|`curandSetQuasiRandomGeneratorDimensions`| | | | | | | | | | |
+|`curandSetQuasiRandomGeneratorDimensions`| | | | |`rocrand_set_quasi_random_generator_dimensions`|1.5.0| | | | |
 |`curandSetStream`| | | | |`rocrand_set_stream`|1.5.0| | | | |
 
 ## **3. Device API Functions**

diff --git a/src/CUDA2HIP_RAND_API_functions.cpp b/src/CUDA2HIP_RAND_API_functions.cpp
@@ -27,8 +27,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RAND_FUNCTION_MAP {
   // RAND Host functions
   {"curandCreateGenerator",                         {"hiprandCreateGenerator",                         "rocrand_create_generator",                                   CONV_LIB_FUNC, API_RAND, 2}},
   {"curandCreateGeneratorHost",                     {"hiprandCreateGeneratorHost",                     "rocrand_create_generator_host_blocking",                     CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandCreatePoissonDistribution",               {"hiprandCreatePoissonDistribution",               "", CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandDestroyDistribution",                     {"hiprandDestroyDistribution",                     "", CONV_LIB_FUNC, API_RAND, 2}},
+  {"curandCreatePoissonDistribution",               {"hiprandCreatePoissonDistribution",               "rocrand_create_poisson_distribution",                        CONV_LIB_FUNC, API_RAND, 2}},
+  {"curandDestroyDistribution",                     {"hiprandDestroyDistribution",                     "rocrand_destroy_discrete_distribution",                      CONV_LIB_FUNC, API_RAND, 2}},
   {"curandDestroyGenerator",                        {"hiprandDestroyGenerator",                        "rocrand_destroy_generator",                                  CONV_LIB_FUNC, API_RAND, 2}},
   {"curandGenerate",                                {"hiprandGenerate",                                "rocrand_generate",                                           CONV_LIB_FUNC, API_RAND, 2}},
   {"curandGenerateLogNormal",                       {"hiprandGenerateLogNormal",                       "rocrand_generate_log_normal",                                CONV_LIB_FUNC, API_RAND, 2}},
@@ -40,19 +40,19 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RAND_FUNCTION_MAP {
   {"curandGenerateSeeds",                           {"hiprandGenerateSeeds",                           "rocrand_initialize_generator",                               CONV_LIB_FUNC, API_RAND, 2}},
   {"curandGenerateUniform",                         {"hiprandGenerateUniform",                         "rocrand_generate_uniform",                                   CONV_LIB_FUNC, API_RAND, 2}},
   {"curandGenerateUniformDouble",                   {"hiprandGenerateUniformDouble",                   "rocrand_generate_uniform_double",                            CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandGetDirectionVectors32",                   {"hiprandGetDirectionVectors32",                   "", CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandGetDirectionVectors64",                   {"hiprandGetDirectionVectors64",                   "", CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandGetProperty",                             {"hiprandGetProperty",                             "", CONV_LIB_FUNC, API_RAND, 2, HIP_UNSUPPORTED}},
-  {"curandGetScrambleConstants32",                  {"hiprandGetScrambleConstants32",                  "", CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandGetScrambleConstants64",                  {"hiprandGetScrambleConstants64",                  "", CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandGetVersion",                              {"hiprandGetVersion",                              "", CONV_LIB_FUNC, API_RAND, 2}},
+  {"curandGetDirectionVectors32",                   {"hiprandGetDirectionVectors32",                   "rocrand_get_direction_vectors32",                            CONV_LIB_FUNC, API_RAND, 2}},
+  {"curandGetDirectionVectors64",                   {"hiprandGetDirectionVectors64",                   "rocrand_get_direction_vectors64",                            CONV_LIB_FUNC, API_RAND, 2}},
+  {"curandGetProperty",                             {"hiprandGetProperty",                             "",                                                           CONV_LIB_FUNC, API_RAND, 2, UNSUPPORTED}},
+  {"curandGetScrambleConstants32",                  {"hiprandGetScrambleConstants32",                  "rocrand_get_scramble_constants32",                           CONV_LIB_FUNC, API_RAND, 2}},
+  {"curandGetScrambleConstants64",                  {"hiprandGetScrambleConstants64",                  "rocrand_get_scramble_constants64",                           CONV_LIB_FUNC, API_RAND, 2}},
+  {"curandGetVersion",                              {"hiprandGetVersion",                              "rocrand_get_version",                                        CONV_LIB_FUNC, API_RAND, 2}},
   {"curandSetGeneratorOffset",                      {"hiprandSetGeneratorOffset",                      "rocrand_set_offset",                                         CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandSetGeneratorOrdering",                    {"hiprandSetGeneratorOrdering",                    "", CONV_LIB_FUNC, API_RAND, 2}},
+  {"curandSetGeneratorOrdering",                    {"hiprandSetGeneratorOrdering",                    "rocrand_set_ordering",                                       CONV_LIB_FUNC, API_RAND, 2}},
   {"curandSetPseudoRandomGeneratorSeed",            {"hiprandSetPseudoRandomGeneratorSeed",            "rocrand_set_seed",                                           CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandSetQuasiRandomGeneratorDimensions",       {"hiprandSetQuasiRandomGeneratorDimensions",       "", CONV_LIB_FUNC, API_RAND, 2}},
+  {"curandSetQuasiRandomGeneratorDimensions",       {"hiprandSetQuasiRandomGeneratorDimensions",       "rocrand_set_quasi_random_generator_dimensions",              CONV_LIB_FUNC, API_RAND, 2}},
   {"curandSetStream",                               {"hiprandSetStream",                               "rocrand_set_stream",                                         CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandMakeMTGP32Constants",                     {"hiprandMakeMTGP32Constants",                     "", CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandMakeMTGP32KernelState",                   {"hiprandMakeMTGP32KernelState",                   "", CONV_LIB_FUNC, API_RAND, 2}},
+  {"curandMakeMTGP32Constants",                     {"hiprandMakeMTGP32Constants",                     "rocrand_make_constant",                                      CONV_LIB_FUNC, API_RAND, 2}},
+  {"curandMakeMTGP32KernelState",                   {"hiprandMakeMTGP32KernelState",                   "rocrand_make_state_mtgp32",                                  CONV_LIB_FUNC, API_RAND, 2}},
 
   // RAND Device functions
   {"curand",                                        {"hiprand",                                        "", CONV_LIB_DEVICE_FUNC, API_RAND, 3}},
@@ -159,6 +159,17 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_RAND_FUNCTION_VER_MAP {
   {"rocrand_set_stream",                            {HIP_1050, HIP_0,    HIP_0,  }},
   {"rocrand_set_seed",                              {HIP_1050, HIP_0,    HIP_0,  }},
   {"rocrand_set_offset",                            {HIP_1050, HIP_0,    HIP_0,  }},
+  {"rocrand_set_ordering",                          {HIP_5050, HIP_0,    HIP_0,  }},
+  {"rocrand_set_quasi_random_generator_dimensions", {HIP_1050, HIP_0,    HIP_0,  }},
+  {"rocrand_get_version",                           {HIP_1050, HIP_0,    HIP_0,  }},
+  {"rocrand_create_poisson_distribution",           {HIP_1050, HIP_0,    HIP_0,  }},
+  {"rocrand_get_direction_vectors32",               {HIP_6000, HIP_0,    HIP_0,  }},
+  {"rocrand_get_direction_vectors64",               {HIP_6000, HIP_0,    HIP_0,  }},
+  {"rocrand_get_scramble_constants32",              {HIP_6000, HIP_0,    HIP_0,  }},
+  {"rocrand_get_scramble_constants64",              {HIP_6000, HIP_0,    HIP_0,  }},
+  {"rocrand_destroy_discrete_distribution",         {HIP_1050, HIP_0,    HIP_0,  }},
+  {"rocrand_make_constant",                         {HIP_1050, HIP_0,    HIP_0,  }},
+  {"rocrand_make_state_mtgp32",                     {HIP_1050, HIP_0,    HIP_0,  }},
 };
 
 const std::map<unsigned int, llvm::StringRef> CUDA_RAND_API_SECTION_MAP {