Merge pull request ROCm#1663 from emankov/HIPIFY

[HIPIFY][rocRAND][feature] Support for `cuRAND -> rocRAND` hipificati…
emankov · Sep 23, 2024 · 79a9a3d · 79a9a3d
2 parents c0c0530 + 585c7e0
commit 79a9a3d
Show file tree

Hide file tree

Showing 8 changed files with 131 additions and 40 deletions.
diff --git a/bin/hipify-perl b/bin/hipify-perl
@@ -2117,6 +2117,12 @@ sub rocSubstitutions {
     subst("curandCreateGeneratorHost", "rocrand_create_generator_host_blocking", "library");
     subst("curandDestroyGenerator", "rocrand_destroy_generator", "library");
     subst("curandGenerate", "rocrand_generate", "library");
+    subst("curandGenerateLogNormal", "rocrand_generate_log_normal", "library");
+    subst("curandGenerateLongLong", "rocrand_generate_long_long", "library");
+    subst("curandGenerateNormal", "rocrand_generate_normal", "library");
+    subst("curandGenerateNormalDouble", "rocrand_generate_normal_double", "library");
+    subst("curandGenerateUniform", "rocrand_generate_uniform", "library");
+    subst("curandGenerateUniformDouble", "rocrand_generate_uniform_double", "library");
     subst("cusolverDnCpotrf", "rocsolver_cpotrf", "library");
     subst("cusolverDnCreate", "rocblas_create_handle", "library");
     subst("cusolverDnDestroy", "rocblas_destroy_handle", "library");
@@ -4839,6 +4845,7 @@ sub simpleSubstitutions {
     subst("curandGenerate", "hiprandGenerate", "library");
     subst("curandGenerateLogNormal", "hiprandGenerateLogNormal", "library");
     subst("curandGenerateLogNormalDouble", "hiprandGenerateLogNormalDouble", "library");
+    subst("curandGenerateLongLong", "hiprandGenerateLongLong", "library");
     subst("curandGenerateNormal", "hiprandGenerateNormal", "library");
     subst("curandGenerateNormalDouble", "hiprandGenerateNormalDouble", "library");
     subst("curandGeneratePoisson", "hiprandGeneratePoisson", "library");
@@ -9313,7 +9320,6 @@ sub warnUnsupportedFunctions {
         "curandHistogramM2K_t",
         "curandHistogramM2K_st",
         "curandGetProperty",
-        "curandGenerateLongLong",
         "curandDistribution_t",
         "curandDistribution_st",
         "curandDistributionShift_t",

diff --git a/docs/tables/CURAND_API_supported_by_HIP.md b/docs/tables/CURAND_API_supported_by_HIP.md
@@ -112,7 +112,7 @@
 |`curandGenerate`| | | | |`hiprandGenerate`|1.5.0| | | | |
 |`curandGenerateLogNormal`| | | | |`hiprandGenerateLogNormal`|1.5.0| | | | |
 |`curandGenerateLogNormalDouble`| | | | |`hiprandGenerateLogNormalDouble`|1.5.0| | | | |
-|`curandGenerateLongLong`| | | | | | | | | | |
+|`curandGenerateLongLong`| | | | |`hiprandGenerateLongLong`|5.5.0| | | | |
 |`curandGenerateNormal`| | | | |`hiprandGenerateNormal`|1.5.0| | | | |
 |`curandGenerateNormalDouble`| | | | |`hiprandGenerateNormalDouble`|1.5.0| | | | |
 |`curandGeneratePoisson`| | | | |`hiprandGeneratePoisson`|1.5.0| | | | |

diff --git a/docs/tables/CURAND_API_supported_by_HIP_and_ROC.md b/docs/tables/CURAND_API_supported_by_HIP_and_ROC.md
@@ -110,15 +110,15 @@
 |`curandDestroyDistribution`| | | | |`hiprandDestroyDistribution`|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| | | | | | | | | | |
+|`curandGenerateLogNormal`| | | | |`hiprandGenerateLogNormal`|1.5.0| | | | |`rocrand_generate_log_normal`|1.5.0| | | | |
 |`curandGenerateLogNormalDouble`| | | | |`hiprandGenerateLogNormalDouble`|1.5.0| | | | | | | | | | |
-|`curandGenerateLongLong`| | | | | | | | | | | | | | | | |
-|`curandGenerateNormal`| | | | |`hiprandGenerateNormal`|1.5.0| | | | | | | | | | |
-|`curandGenerateNormalDouble`| | | | |`hiprandGenerateNormalDouble`|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| | | | | | | | | | |
-|`curandGenerateUniform`| | | | |`hiprandGenerateUniform`|1.5.0| | | | | | | | | | |
-|`curandGenerateUniformDouble`| | | | |`hiprandGenerateUniformDouble`|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| | | | | | | | | | |
 |`curandGetProperty`|8.0| | | | | | | | | | | | | | | |

diff --git a/docs/tables/CURAND_API_supported_by_ROC.md b/docs/tables/CURAND_API_supported_by_ROC.md
@@ -110,15 +110,15 @@
 |`curandDestroyDistribution`| | | | | | | | | | |
 |`curandDestroyGenerator`| | | | |`rocrand_destroy_generator`|1.5.0| | | | |
 |`curandGenerate`| | | | |`rocrand_generate`|1.5.0| | | | |
-|`curandGenerateLogNormal`| | | | | | | | | | |
+|`curandGenerateLogNormal`| | | | |`rocrand_generate_log_normal`|1.5.0| | | | |
 |`curandGenerateLogNormalDouble`| | | | | | | | | | |
-|`curandGenerateLongLong`| | | | | | | | | | |
-|`curandGenerateNormal`| | | | | | | | | | |
-|`curandGenerateNormalDouble`| | | | | | | | | | |
+|`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`| | | | | | | | | | |
-|`curandGenerateUniform`| | | | | | | | | | |
-|`curandGenerateUniformDouble`| | | | | | | | | | |
+|`curandGenerateUniform`| | | | |`rocrand_generate_uniform`|1.5.0| | | | |
+|`curandGenerateUniformDouble`| | | | |`rocrand_generate_uniform_double`|1.5.0| | | | |
 |`curandGetDirectionVectors32`| | | | | | | | | | |
 |`curandGetDirectionVectors64`| | | | | | | | | | |
 |`curandGetProperty`|8.0| | | | | | | | | |

diff --git a/src/CUDA2HIP_RAND_API_functions.cpp b/src/CUDA2HIP_RAND_API_functions.cpp
@@ -31,15 +31,15 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RAND_FUNCTION_MAP {
   {"curandDestroyDistribution",                     {"hiprandDestroyDistribution",                     "", 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",                       "", 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}},
-  {"curandGenerateLongLong",                        {"hiprandGenerateLongLong",                        "", CONV_LIB_FUNC, API_RAND, 2, HIP_UNSUPPORTED}},
-  {"curandGenerateNormal",                          {"hiprandGenerateNormal",                          "", CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandGenerateNormalDouble",                    {"hiprandGenerateNormalDouble",                    "", 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}},
-  {"curandGenerateUniform",                         {"hiprandGenerateUniform",                         "", CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandGenerateUniformDouble",                   {"hiprandGenerateUniformDouble",                   "", 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}},
@@ -141,11 +141,18 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_RAND_FUNCTION_VER_MAP {
   {"hiprandGetScrambleConstants32",                 {HIP_6000, HIP_0,    HIP_0   }},
   {"hiprandGetScrambleConstants64",                 {HIP_6000, HIP_0,    HIP_0   }},
   {"hiprandSetGeneratorOrdering",                   {HIP_6020, HIP_0,    HIP_0,  }},
+  {"hiprandGenerateLongLong",                       {HIP_5050, HIP_0,    HIP_0,  }},
 
   {"rocrand_create_generator",                      {HIP_1050, HIP_0,    HIP_0,  }},
   {"rocrand_create_generator_host_blocking",        {HIP_6020, HIP_0,    HIP_0,  }},
   {"rocrand_destroy_generator",                     {HIP_1050, HIP_0,    HIP_0,  }},
   {"rocrand_generate",                              {HIP_1050, HIP_0,    HIP_0,  }},
+  {"rocrand_generate_long_long",                    {HIP_5040, HIP_0,    HIP_0,  }},
+  {"rocrand_generate_uniform",                      {HIP_1050, HIP_0,    HIP_0,  }},
+  {"rocrand_generate_uniform_double",               {HIP_1050, HIP_0,    HIP_0,  }},
+  {"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,  }},
 };
 
 const std::map<unsigned int, llvm::StringRef> CUDA_RAND_API_SECTION_MAP {

diff --git a/tests/unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp b/tests/unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp
@@ -74,13 +74,13 @@ void run_benchmark(const cli::Parser& parser, const rng_type_t rng_type,
     CUDA_CALL(cudaMalloc((void**)&data, size * sizeof(T)));
 
     // CHECK: hiprandGenerator_t generator;
-    // CHECK: CURAND_CALL(hiprandCreateGenerator(&generator, rng_type));
+    // CHECK-NEXT: CURAND_CALL(hiprandCreateGenerator(&generator, rng_type));
     curandGenerator_t generator;
     CURAND_CALL(curandCreateGenerator(&generator, rng_type));
 
     const size_t dimensions = parser.get<size_t>("dimensions");
     // CHECK: hiprandStatus_t status = hiprandSetQuasiRandomGeneratorDimensions(generator, dimensions);
-    // CHECK: if (status != HIPRAND_STATUS_TYPE_ERROR)
+    // CHECK-NEXT: if (status != HIPRAND_STATUS_TYPE_ERROR)
     curandStatus_t status = curandSetQuasiRandomGeneratorDimensions(generator, dimensions);
     if (status != CURAND_STATUS_TYPE_ERROR)  // If the RNG is not quasi-random
     {
@@ -113,7 +113,7 @@ void run_benchmark(const cli::Parser& parser, const rng_type_t rng_type,
               << " ms, Time (all) = " << std::setw(8) << elapsed.count() << " ms, Size = " << size
               << std::endl;
     // CHECK: CURAND_CALL(hiprandDestroyGenerator(generator));
-    // CHECK: CUDA_CALL(hipFree(data));
+    // CHECK-NEXT: CUDA_CALL(hipFree(data));
     CURAND_CALL(curandDestroyGenerator(generator));
     CUDA_CALL(cudaFree(data));
 }
@@ -122,7 +122,7 @@ void run_benchmarks(const cli::Parser& parser, const rng_type_t rng_type,
                     const std::string& distribution) {
     if (distribution == "uniform-uint") {
         // CHECK: if (rng_type != HIPRAND_RNG_QUASI_SOBOL64 &&
-        // CHECK: rng_type != HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64) {
+        // CHECK-NEXT: rng_type != HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64) {
         if (rng_type != CURAND_RNG_QUASI_SOBOL64 &&
             rng_type != CURAND_RNG_QUASI_SCRAMBLED_SOBOL64) {
             run_benchmark<unsigned int>(
@@ -136,39 +136,38 @@ void run_benchmarks(const cli::Parser& parser, const rng_type_t rng_type,
     }
     if (distribution == "uniform-long-long") {
         // CHECK: if (rng_type == HIPRAND_RNG_QUASI_SOBOL64 ||
-        // CHECK: rng_type == HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64)
+        // CHECK-NEXT: rng_type == HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64)
         if (rng_type == CURAND_RNG_QUASI_SOBOL64 ||
             rng_type == CURAND_RNG_QUASI_SCRAMBLED_SOBOL64) {
             run_benchmark<unsigned long long>(
                 parser, rng_type,
                 // CHECK: [](hiprandGenerator_t gen, unsigned long long* data, size_t size) {
                 [](curandGenerator_t gen, unsigned long long* data, size_t size) {
-                    // curandGenerateLongLong is yet unsupported by HIP
-                    // CHECK-NOT: return hiprandGenerateLongLong(gen, data, size);
+                    // CHECK: return hiprandGenerateLongLong(gen, data, size);
                     return curandGenerateLongLong(gen, data, size);
                 });
         }
     }
     if (distribution == "uniform-float") {
         run_benchmark<float>(parser, rng_type,
                              // CHECK: [](hiprandGenerator_t gen, float* data, size_t size) {
-                             // CHECK: return hiprandGenerateUniform(gen, data, size);
+                             // CHECK-NEXT: return hiprandGenerateUniform(gen, data, size);
                              [](curandGenerator_t gen, float* data, size_t size) {
                                  return curandGenerateUniform(gen, data, size);
                              });
     }
     if (distribution == "uniform-double") {
         run_benchmark<double>(parser, rng_type,
                               // CHECK: [](hiprandGenerator_t gen, double* data, size_t size) {
-                              // CHECK: return hiprandGenerateUniformDouble(gen, data, size);
+                              // CHECK-NEXT: return hiprandGenerateUniformDouble(gen, data, size);
                               [](curandGenerator_t gen, double* data, size_t size) {
                                   return curandGenerateUniformDouble(gen, data, size);
                               });
     }
     if (distribution == "normal-float") {
         run_benchmark<float>(parser, rng_type,
                              // CHECK: [](hiprandGenerator_t gen, float* data, size_t size) {
-                             // CHECK: return hiprandGenerateNormal(gen, data, size, 0.0f, 1.0f);
+                             // CHECK-NEXT: return hiprandGenerateNormal(gen, data, size, 0.0f, 1.0f);
                              [](curandGenerator_t gen, float* data, size_t size) {
                                  return curandGenerateNormal(gen, data, size, 0.0f, 1.0f);
                              });
@@ -177,15 +176,15 @@ void run_benchmarks(const cli::Parser& parser, const rng_type_t rng_type,
         run_benchmark<double>(
             parser, rng_type,
             // CHECK: [](hiprandGenerator_t gen, double* data, size_t size) {
-            // CHECK: return hiprandGenerateNormalDouble(gen, data, size, 0.0, 1.0);
+            // CHECK-NEXT: return hiprandGenerateNormalDouble(gen, data, size, 0.0, 1.0);
             [](curandGenerator_t gen, double* data, size_t size) {
                 return curandGenerateNormalDouble(gen, data, size, 0.0, 1.0);
             });
     }
     if (distribution == "log-normal-float") {
         run_benchmark<float>(parser, rng_type,
                              // CHECK: [](hiprandGenerator_t gen, float* data, size_t size) {
-                             // CHECK: return hiprandGenerateLogNormal(gen, data, size, 0.0f, 1.0f);
+                             // CHECK-NEXT: return hiprandGenerateLogNormal(gen, data, size, 0.0f, 1.0f);
                              [](curandGenerator_t gen, float* data, size_t size) {
                                  return curandGenerateLogNormal(gen, data, size, 0.0f, 1.0f);
                              });
@@ -194,7 +193,7 @@ void run_benchmarks(const cli::Parser& parser, const rng_type_t rng_type,
         run_benchmark<double>(
             parser, rng_type,
             // CHECK: [](hiprandGenerator_t gen, double* data, size_t size) {
-            // CHECK: return hiprandGenerateLogNormalDouble(gen, data, size, 0.0, 1.0);
+            // CHECK-NEXT: return hiprandGenerateLogNormalDouble(gen, data, size, 0.0, 1.0);
             [](curandGenerator_t gen, double* data, size_t size) {
                 return curandGenerateLogNormalDouble(gen, data, size, 0.0, 1.0);
             });
@@ -207,7 +206,7 @@ void run_benchmarks(const cli::Parser& parser, const rng_type_t rng_type,
             run_benchmark<unsigned int>(
                 parser, rng_type,
                 // CHECK: [lambda](hiprandGenerator_t gen, unsigned int* data, size_t size) {
-                // CHECK: return hiprandGeneratePoisson(gen, data, size, lambda);
+                // CHECK-NEXT: return hiprandGeneratePoisson(gen, data, size, lambda);
                 [lambda](curandGenerator_t gen, unsigned int* data, size_t size) {
                     return curandGeneratePoisson(gen, data, size, lambda);
                 });
@@ -219,9 +218,9 @@ const std::vector<std::string> all_engines = {
     "xorwow", "mrg32k3a", "mtgp32",
     // "mt19937",
     "philox", "sobol32",
-    // "scrambled_sobol32",
-    // "sobol64",
-    // "scrambled_sobol64",
+    "scrambled_sobol32",
+    "sobol64",
+    "scrambled_sobol64",
 };
 
 const std::vector<std::string> all_distributions = {
@@ -287,8 +286,8 @@ int main(int argc, char* argv[]) {
     CUDA_CALL(cudaRuntimeGetVersion(&runtime_version));
     int device_id;
     // CHECK: CUDA_CALL(hipGetDevice(&device_id));
-    // CHECK: hipDeviceProp_t props;
-    // CHECK: CUDA_CALL(hipGetDeviceProperties(&props, device_id));
+    // CHECK-NEXT: hipDeviceProp_t props;
+    // CHECK-NEXT: CUDA_CALL(hipGetDeviceProperties(&props, device_id));
     CUDA_CALL(cudaGetDevice(&device_id));
     cudaDeviceProp props;
     CUDA_CALL(cudaGetDeviceProperties(&props, device_id));