Merge pull request #1671 from emankov/HIPIFY

[HIPIFY][rocRAND][feature] Support for `cuRAND -> rocRAND` hipification - Step 9 - Functions
ROCm · Sep 27, 2024 · 10b95fb · 10b95fb
2 parents 789ef05 + 3d3818e
commit 10b95fb
Show file tree

Hide file tree

Showing 6 changed files with 108 additions and 32 deletions.
diff --git a/bin/hipify-perl b/bin/hipify-perl
@@ -2136,11 +2136,17 @@ sub rocSubstitutions {
     subst("curandDestroyGenerator", "rocrand_destroy_generator", "library");
     subst("curandGenerate", "rocrand_generate", "library");
     subst("curandGenerateLogNormal", "rocrand_generate_log_normal", "library");
+    subst("curandGenerateLogNormalDouble", "rocrand_generate_log_normal_double", "library");
     subst("curandGenerateLongLong", "rocrand_generate_long_long", "library");
     subst("curandGenerateNormal", "rocrand_generate_normal", "library");
     subst("curandGenerateNormalDouble", "rocrand_generate_normal_double", "library");
+    subst("curandGeneratePoisson", "rocrand_generate_poisson", "library");
+    subst("curandGenerateSeeds", "rocrand_initialize_generator", "library");
     subst("curandGenerateUniform", "rocrand_generate_uniform", "library");
     subst("curandGenerateUniformDouble", "rocrand_generate_uniform_double", "library");
+    subst("curandSetGeneratorOffset", "rocrand_set_offset", "library");
+    subst("curandSetPseudoRandomGeneratorSeed", "rocrand_set_seed", "library");
+    subst("curandSetStream", "rocrand_set_stream", "library");
     subst("cusolverDnCpotrf", "rocsolver_cpotrf", "library");
     subst("cusolverDnCreate", "rocblas_create_handle", "library");
     subst("cusolverDnDestroy", "rocblas_destroy_handle", "library");

diff --git a/docs/tables/CURAND_API_supported_by_HIP_and_ROC.md b/docs/tables/CURAND_API_supported_by_HIP_and_ROC.md
@@ -111,12 +111,12 @@
 |`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| | | | |
-|`curandGenerateLogNormalDouble`| | | | |`hiprandGenerateLogNormalDouble`|1.5.0| | | | | | | | | | |
+|`curandGenerateLogNormalDouble`| | | | |`hiprandGenerateLogNormalDouble`|1.5.0| | | | |`rocrand_generate_log_normal_double`|1.5.0| | | | |
 |`curandGenerateLongLong`| | | | |`hiprandGenerateLongLong`|5.5.0| | | | |`rocrand_generate_long_long`|5.4.0| | | | |
 |`curandGenerateNormal`| | | | |`hiprandGenerateNormal`|1.5.0| | | | |`rocrand_generate_normal`|1.5.0| | | | |
 |`curandGenerateNormalDouble`| | | | |`hiprandGenerateNormalDouble`|1.5.0| | | | |`rocrand_generate_normal_double`|1.5.0| | | | |
-|`curandGeneratePoisson`| | | | |`hiprandGeneratePoisson`|1.5.0| | | | | | | | | | |
-|`curandGenerateSeeds`| | | | |`hiprandGenerateSeeds`|1.5.0| | | | | | | | | | |
+|`curandGeneratePoisson`| | | | |`hiprandGeneratePoisson`|1.5.0| | | | |`rocrand_generate_poisson`|1.5.0| | | | |
+|`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| | | | | | | | | | |
@@ -127,11 +127,11 @@
 |`curandGetVersion`| | | | |`hiprandGetVersion`|1.5.0| | | | | | | | | | |
 |`curandMakeMTGP32Constants`| | | | |`hiprandMakeMTGP32Constants`|1.5.0| | | | | | | | | | |
 |`curandMakeMTGP32KernelState`| | | | |`hiprandMakeMTGP32KernelState`|1.5.0| | | | | | | | | | |
-|`curandSetGeneratorOffset`| | | | |`hiprandSetGeneratorOffset`|1.5.0| | | | | | | | | | |
+|`curandSetGeneratorOffset`| | | | |`hiprandSetGeneratorOffset`|1.5.0| | | | |`rocrand_set_offset`|1.5.0| | | | |
 |`curandSetGeneratorOrdering`| | | | |`hiprandSetGeneratorOrdering`|6.2.0| | | | | | | | | | |
-|`curandSetPseudoRandomGeneratorSeed`| | | | |`hiprandSetPseudoRandomGeneratorSeed`|1.5.0| | | | | | | | | | |
+|`curandSetPseudoRandomGeneratorSeed`| | | | |`hiprandSetPseudoRandomGeneratorSeed`|1.5.0| | | | |`rocrand_set_seed`|1.5.0| | | | |
 |`curandSetQuasiRandomGeneratorDimensions`| | | | |`hiprandSetQuasiRandomGeneratorDimensions`|1.5.0| | | | | | | | | | |
-|`curandSetStream`| | | | |`hiprandSetStream`|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
@@ -111,12 +111,12 @@
 |`curandDestroyGenerator`| | | | |`rocrand_destroy_generator`|1.5.0| | | | |
 |`curandGenerate`| | | | |`rocrand_generate`|1.5.0| | | | |
 |`curandGenerateLogNormal`| | | | |`rocrand_generate_log_normal`|1.5.0| | | | |
-|`curandGenerateLogNormalDouble`| | | | | | | | | | |
+|`curandGenerateLogNormalDouble`| | | | |`rocrand_generate_log_normal_double`|1.5.0| | | | |
 |`curandGenerateLongLong`| | | | |`rocrand_generate_long_long`|5.4.0| | | | |
 |`curandGenerateNormal`| | | | |`rocrand_generate_normal`|1.5.0| | | | |
 |`curandGenerateNormalDouble`| | | | |`rocrand_generate_normal_double`|1.5.0| | | | |
-|`curandGeneratePoisson`| | | | | | | | | | |
-|`curandGenerateSeeds`| | | | | | | | | | |
+|`curandGeneratePoisson`| | | | |`rocrand_generate_poisson`|1.5.0| | | | |
+|`curandGenerateSeeds`| | | | |`rocrand_initialize_generator`|1.5.0| | | | |
 |`curandGenerateUniform`| | | | |`rocrand_generate_uniform`|1.5.0| | | | |
 |`curandGenerateUniformDouble`| | | | |`rocrand_generate_uniform_double`|1.5.0| | | | |
 |`curandGetDirectionVectors32`| | | | | | | | | | |
@@ -127,11 +127,11 @@
 |`curandGetVersion`| | | | | | | | | | |
 |`curandMakeMTGP32Constants`| | | | | | | | | | |
 |`curandMakeMTGP32KernelState`| | | | | | | | | | |
-|`curandSetGeneratorOffset`| | | | | | | | | | |
+|`curandSetGeneratorOffset`| | | | |`rocrand_set_offset`|1.5.0| | | | |
 |`curandSetGeneratorOrdering`| | | | | | | | | | |
-|`curandSetPseudoRandomGeneratorSeed`| | | | | | | | | | |
+|`curandSetPseudoRandomGeneratorSeed`| | | | |`rocrand_set_seed`|1.5.0| | | | |
 |`curandSetQuasiRandomGeneratorDimensions`| | | | | | | | | | |
-|`curandSetStream`| | | | | | | | | | |
+|`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
@@ -32,12 +32,12 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RAND_FUNCTION_MAP {
   {"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}},
-  {"curandGenerateLogNormalDouble",                 {"hiprandGenerateLogNormalDouble",                 "", CONV_LIB_FUNC, API_RAND, 2}},
+  {"curandGenerateLogNormalDouble",                 {"hiprandGenerateLogNormalDouble",                 "rocrand_generate_log_normal_double",                         CONV_LIB_FUNC, API_RAND, 2}},
   {"curandGenerateLongLong",                        {"hiprandGenerateLongLong",                        "rocrand_generate_long_long",                                 CONV_LIB_FUNC, API_RAND, 2}},
   {"curandGenerateNormal",                          {"hiprandGenerateNormal",                          "rocrand_generate_normal",                                    CONV_LIB_FUNC, API_RAND, 2}},
   {"curandGenerateNormalDouble",                    {"hiprandGenerateNormalDouble",                    "rocrand_generate_normal_double",                             CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandGeneratePoisson",                         {"hiprandGeneratePoisson",                         "", CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandGenerateSeeds",                           {"hiprandGenerateSeeds",                           "", CONV_LIB_FUNC, API_RAND, 2}},
+  {"curandGeneratePoisson",                         {"hiprandGeneratePoisson",                         "rocrand_generate_poisson",                                   CONV_LIB_FUNC, API_RAND, 2}},
+  {"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}},
@@ -46,11 +46,11 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RAND_FUNCTION_MAP {
   {"curandGetScrambleConstants32",                  {"hiprandGetScrambleConstants32",                  "", CONV_LIB_FUNC, API_RAND, 2}},
   {"curandGetScrambleConstants64",                  {"hiprandGetScrambleConstants64",                  "", CONV_LIB_FUNC, API_RAND, 2}},
   {"curandGetVersion",                              {"hiprandGetVersion",                              "", CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandSetGeneratorOffset",                      {"hiprandSetGeneratorOffset",                      "", 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}},
-  {"curandSetPseudoRandomGeneratorSeed",            {"hiprandSetPseudoRandomGeneratorSeed",            "", 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}},
-  {"curandSetStream",                               {"hiprandSetStream",                               "", 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}},
 
@@ -153,6 +153,12 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_RAND_FUNCTION_VER_MAP {
   {"rocrand_generate_normal",                       {HIP_1050, HIP_0,    HIP_0,  }},
   {"rocrand_generate_normal_double",                {HIP_1050, HIP_0,    HIP_0,  }},
   {"rocrand_generate_log_normal",                   {HIP_1050, HIP_0,    HIP_0,  }},
+  {"rocrand_generate_log_normal_double",            {HIP_1050, HIP_0,    HIP_0,  }},
+  {"rocrand_generate_poisson",                      {HIP_1050, HIP_0,    HIP_0,  }},
+  {"rocrand_initialize_generator",                  {HIP_1050, HIP_0,    HIP_0,  }},
+  {"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,  }},
 };
 
 const std::map<unsigned int, llvm::StringRef> CUDA_RAND_API_SECTION_MAP {

diff --git a/tests/unit_tests/synthetic/libraries/curand2hiprand.cu b/tests/unit_tests/synthetic/libraries/curand2hiprand.cu
@@ -21,15 +21,17 @@ int main() {
   unsigned int *outputPtr = nullptr;
   float *outputPtrFloat = nullptr;
   double *outputPtrDouble = nullptr;
+  unsigned long long *outputPtrUll = nullptr;
+  unsigned long long offset = 0;
   size_t num = 0;
   float mean = 0.f;
+  double dmean = 0.f;
   float stddev = 0.f;
+  double dstddev = 0.f;
+  double dlambda = 0.f;
 
-#if defined(_WIN32)
-  unsigned long long *outputPtrUll = nullptr;
-#else
-  unsigned long *outputPtrUll = nullptr;
-#endif
+  // CHECK: hipStream_t stream;
+  cudaStream_t stream;
 
   // CHECK: hiprandStatus randStatus;
   // CHECK-NEXT: hiprandStatus_t status;
@@ -191,8 +193,8 @@ int main() {
 
   // CUDA: curandStatus_t CURANDAPI curandGenerateNormalDouble(curandGenerator_t generator, double *outputPtr, size_t n, double mean, double stddev);
   // HIP: hiprandStatus_t HIPRANDAPI hiprandGenerateNormalDouble(hiprandGenerator_t generator, double * output_data, size_t n, double mean, double stddev);
-  // CHECK: status = hiprandGenerateNormalDouble(randGenerator, outputPtrDouble, num, mean, stddev);
-  status = curandGenerateNormalDouble(randGenerator, outputPtrDouble, num, mean, stddev);
+  // CHECK: status = hiprandGenerateNormalDouble(randGenerator, outputPtrDouble, num, dmean, dstddev);
+  status = curandGenerateNormalDouble(randGenerator, outputPtrDouble, num, dmean, dstddev);
 
   // CUDA: curandStatus_t CURANDAPI curandGenerateUniform(curandGenerator_t generator, float *outputPtr, size_t num);
   // HIP: hiprandStatus_t HIPRANDAPI hiprandGenerateUniform(hiprandGenerator_t generator, float * output_data, size_t n);
@@ -204,6 +206,36 @@ int main() {
   // CHECK: status = hiprandGenerateUniformDouble(randGenerator, outputPtrDouble, num);
   status = curandGenerateUniformDouble(randGenerator, outputPtrDouble, num);
 
+  // CUDA: curandStatus_t CURANDAPI curandGenerateLogNormalDouble(curandGenerator_t generator, double *outputPtr, size_t n, double mean, double stddev);
+  // HIP: hiprandStatus_t HIPRANDAPI hiprandGenerateLogNormalDouble(hiprandGenerator_t generator, double * output_data, size_t n, double mean, double stddev);
+  // CHECK: status = hiprandGenerateLogNormalDouble(randGenerator, outputPtrDouble, num, dmean, dstddev);
+  status = curandGenerateLogNormalDouble(randGenerator, outputPtrDouble, num, dmean, dstddev);
+
+  // CUDA: curandStatus_t CURANDAPI curandGeneratePoisson(curandGenerator_t generator, unsigned int *outputPtr, size_t n, double lambda);
+  // HIP: hiprandStatus_t HIPRANDAPI hiprandGeneratePoisson(hiprandGenerator_t generator, unsigned int * output_data, size_t n, double lambda);
+  // CHECK: status = hiprandGeneratePoisson(randGenerator, outputPtr, num, dlambda);
+  status = curandGeneratePoisson(randGenerator, outputPtr, num, dlambda);
+
+  // CUDA: curandStatus_t CURANDAPI curandGenerateSeeds(curandGenerator_t generator);
+  // HIP: hiprandStatus_t HIPRANDAPI hiprandGenerateSeeds(hiprandGenerator_t generator);
+  // CHECK: status = hiprandGenerateSeeds(randGenerator);
+  status = curandGenerateSeeds(randGenerator);
+
+  // CUDA: curandStatus_t CURANDAPI curandSetGeneratorOffset(curandGenerator_t generator, unsigned long long offset);
+  // HIP: hiprandStatus_t HIPRANDAPI hiprandSetGeneratorOffset(hiprandGenerator_t generator, unsigned long long offset);
+  // CHECK: status = hiprandSetGeneratorOffset(randGenerator, offset);
+  status = curandSetGeneratorOffset(randGenerator, offset);
+
+  // CUDA: curandStatus_t CURANDAPI curandSetPseudoRandomGeneratorSeed(curandGenerator_t generator, unsigned long long seed);
+  // HIP: hiprandStatus_t HIPRANDAPI hiprandSetPseudoRandomGeneratorSeed(hiprandGenerator_t generator, unsigned long long seed);
+  // CHECK: status = hiprandSetPseudoRandomGeneratorSeed(randGenerator, offset);
+  status = curandSetPseudoRandomGeneratorSeed(randGenerator, offset);
+
+  // CUDA: curandStatus_t CURANDAPI curandSetStream(curandGenerator_t generator, cudaStream_t stream);
+  // HIP: hiprandStatus_t HIPRANDAPI hiprandSetStream(hiprandGenerator_t generator, hipStream_t stream);
+  // CHECK: status = hiprandSetStream(randGenerator, stream);
+  status = curandSetStream(randGenerator, stream);
+
 #if CUDA_VERSION >= 11000 && CURAND_VERSION >= 10200
   // CHECK: hiprandOrdering_t RAND_ORDERING_PSEUDO_LEGACY = HIPRAND_ORDERING_PSEUDO_LEGACY;
   curandOrdering_t RAND_ORDERING_PSEUDO_LEGACY = CURAND_ORDERING_PSEUDO_LEGACY;