Merge pull request ROCm#1651 from emankov/HIPIFY

[HIPIFY][rocRAND][feature] Support for `cuRAND -> rocRAND` hipification - Step 7 - Functions
emankov · Sep 16, 2024 · 8aecc4b · 8aecc4b
2 parents ce15684 + 187080b
commit 8aecc4b
Show file tree

Hide file tree

Showing 6 changed files with 69 additions and 14 deletions.
diff --git a/bin/hipify-perl b/bin/hipify-perl
@@ -2598,6 +2598,10 @@ sub rocSubstitutions {
     subst("cudnnSoftmaxBackward", "miopenSoftmaxBackward_V2", "library");
     subst("cudnnSoftmaxForward", "miopenSoftmaxForward_V2", "library");
     subst("cudnnTransformTensor", "miopenTransformTensor", "library");
+    subst("curandCreateGenerator", "rocrand_create_generator", "library");
+    subst("curandCreateGeneratorHost", "rocrand_create_generator_host_blocking", "library");
+    subst("curandDestroyGenerator", "rocrand_destroy_generator", "library");
+    subst("curandGenerate", "rocrand_generate", "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
@@ -104,12 +104,12 @@
 
 |**CUDA**|**A**|**D**|**C**|**R**|**HIP**|**A**|**D**|**C**|**R**|**E**|**ROC**|**A**|**D**|**C**|**R**|**E**|
 |:--|:-:|:-:|:-:|:-:|:--|:-:|:-:|:-:|:-:|:-:|:--|:-:|:-:|:-:|:-:|:-:|
-|`curandCreateGenerator`| | | | |`hiprandCreateGenerator`|1.5.0| | | | | | | | | | |
-|`curandCreateGeneratorHost`| | | | |`hiprandCreateGeneratorHost`|1.5.0| | | | | | | | | | |
+|`curandCreateGenerator`| | | | |`hiprandCreateGenerator`|1.5.0| | | | |`rocrand_create_generator`|1.5.0| | | | |
+|`curandCreateGeneratorHost`| | | | |`hiprandCreateGeneratorHost`|1.5.0| | | | |`rocrand_create_generator_host_blocking`|6.2.0| | | |6.2.0|
 |`curandCreatePoissonDistribution`| | | | |`hiprandCreatePoissonDistribution`|1.5.0| | | | | | | | | | |
 |`curandDestroyDistribution`| | | | |`hiprandDestroyDistribution`|1.5.0| | | | | | | | | | |
-|`curandDestroyGenerator`| | | | |`hiprandDestroyGenerator`|1.5.0| | | | | | | | | | |
-|`curandGenerate`| | | | |`hiprandGenerate`|1.5.0| | | | | | | | | | |
+|`curandDestroyGenerator`| | | | |`hiprandDestroyGenerator`|1.5.0| | | | |`rocrand_destroy_generator`|1.5.0| | | |6.2.0|
+|`curandGenerate`| | | | |`hiprandGenerate`|1.5.0| | | | |`rocrand_generate`|1.5.0| | | |6.2.0|
 |`curandGenerateLogNormal`| | | | |`hiprandGenerateLogNormal`|1.5.0| | | | | | | | | | |
 |`curandGenerateLogNormalDouble`| | | | |`hiprandGenerateLogNormalDouble`|1.5.0| | | | | | | | | | |
 |`curandGenerateLongLong`| | | | | | | | | | | | | | | | |

diff --git a/docs/tables/CURAND_API_supported_by_ROC.md b/docs/tables/CURAND_API_supported_by_ROC.md
@@ -104,12 +104,12 @@
 
 |**CUDA**|**A**|**D**|**C**|**R**|**ROC**|**A**|**D**|**C**|**R**|**E**|
 |:--|:-:|:-:|:-:|:-:|:--|:-:|:-:|:-:|:-:|:-:|
-|`curandCreateGenerator`| | | | | | | | | | |
-|`curandCreateGeneratorHost`| | | | | | | | | | |
+|`curandCreateGenerator`| | | | |`rocrand_create_generator`|1.5.0| | | | |
+|`curandCreateGeneratorHost`| | | | |`rocrand_create_generator_host_blocking`|6.2.0| | | |6.2.0|
 |`curandCreatePoissonDistribution`| | | | | | | | | | |
 |`curandDestroyDistribution`| | | | | | | | | | |
-|`curandDestroyGenerator`| | | | | | | | | | |
-|`curandGenerate`| | | | | | | | | | |
+|`curandDestroyGenerator`| | | | |`rocrand_destroy_generator`|1.5.0| | | |6.2.0|
+|`curandGenerate`| | | | |`rocrand_generate`|1.5.0| | | |6.2.0|
 |`curandGenerateLogNormal`| | | | | | | | | | |
 |`curandGenerateLogNormalDouble`| | | | | | | | | | |
 |`curandGenerateLongLong`| | | | | | | | | | |

diff --git a/src/CUDA2HIP_RAND_API_functions.cpp b/src/CUDA2HIP_RAND_API_functions.cpp
@@ -25,12 +25,12 @@ THE SOFTWARE.
 // Map of all functions
 const std::map<llvm::StringRef, hipCounter> CUDA_RAND_FUNCTION_MAP {
   // RAND Host functions
-  {"curandCreateGenerator",                         {"hiprandCreateGenerator",                         "", CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandCreateGeneratorHost",                     {"hiprandCreateGeneratorHost",                     "", CONV_LIB_FUNC, API_RAND, 2}},
+  {"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}},
-  {"curandDestroyGenerator",                        {"hiprandDestroyGenerator",                        "", CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandGenerate",                                {"hiprandGenerate",                                "", 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}},
   {"curandGenerateLogNormalDouble",                 {"hiprandGenerateLogNormalDouble",                 "", CONV_LIB_FUNC, API_RAND, 2}},
   {"curandGenerateLongLong",                        {"hiprandGenerateLongLong",                        "", CONV_LIB_FUNC, API_RAND, 2, HIP_UNSUPPORTED}},
@@ -141,6 +141,11 @@ 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,  HIP_LATEST}},
+
+  {"rocrand_create_generator",                      {HIP_1050, HIP_0,    HIP_0,  }},
+  {"rocrand_create_generator_host_blocking",        {HIP_6020, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"rocrand_destroy_generator",                     {HIP_1050, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"rocrand_generate",                              {HIP_1050, HIP_0,    HIP_0,  HIP_LATEST}},
 };
 
 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
@@ -18,6 +18,9 @@
 int main() {
   printf("21. cuRAND API to hipRAND API synthetic test\n");
 
+  unsigned int *outputPtr = nullptr;
+  size_t num = 0;
+
   // CHECK: hiprandStatus randStatus;
   // CHECK-NEXT: hiprandStatus_t status;
   // CHECK-NEXT: hiprandStatus_t STATUS_SUCCESS = HIPRAND_STATUS_SUCCESS;
@@ -131,6 +134,21 @@ int main() {
   curandDirectionVectorSet_t DIRECTION_VECTORS_64_JOEKUO6 = CURAND_DIRECTION_VECTORS_64_JOEKUO6;
   curandDirectionVectorSet_t SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6 = CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6;
 
+  // CUDA: curandStatus_t CURANDAPI curandCreateGenerator(curandGenerator_t *generator, curandRngType_t rng_type);
+  // HIP: hiprandStatus_t HIPRANDAPI hiprandCreateGenerator(hiprandGenerator_t* generator, hiprandRngType_t rng_type)
+  // CHECK: status = hiprandCreateGenerator(&randGenerator, randRngType_t);
+  status = curandCreateGenerator(&randGenerator, randRngType_t);
+
+  // CUDA: curandStatus_t CURANDAPI curandDestroyGenerator(curandGenerator_t generator);
+  // HIP: hiprandStatus_t HIPRANDAPI hiprandDestroyGenerator(hiprandGenerator_t generator);
+  // CHECK: status = hiprandDestroyGenerator(randGenerator);
+  status = curandDestroyGenerator(randGenerator);
+
+  // CUDA: curandStatus_t CURANDAPI curandCreateGeneratorHost(curandGenerator_t *generator, curandRngType_t rng_type);
+  // HIP: hiprandStatus_t HIPRANDAPI hiprandCreateGeneratorHost(hiprandGenerator_t * generator, hiprandRngType_t rng_type);
+  // CHECK: status = hiprandCreateGeneratorHost(&randGenerator, randRngType_t);
+  status = curandCreateGeneratorHost(&randGenerator, randRngType_t);
+
   // CUDA: curandStatus_t CURANDAPI curandSetGeneratorOrdering(curandGenerator_t generator, curandOrdering_t order);
   // HIP: hiprandStatus_t HIPRANDAPI hiprandSetGeneratorOrdering(hiprandGenerator_t generator, hiprandOrdering_t order);
   // CHECK: status = hiprandSetGeneratorOrdering(randGenerator, randOrdering_t);
@@ -141,6 +159,11 @@ int main() {
   // CHECK: status = hiprandGetDirectionVectors64(&pDirections64, directionVectorSet_t);
   status = curandGetDirectionVectors64(&pDirections64, directionVectorSet_t);
 
+  // CUDA: curandStatus_t CURANDAPI curandGenerate(curandGenerator_t generator, unsigned int *outputPtr, size_t num);
+  // HIP: hiprandStatus_t HIPRANDAPI hiprandGenerate(hiprandGenerator_t generator, unsigned int * output_data, size_t n);
+  // CHECK: status = hiprandGenerate(randGenerator, outputPtr, num);
+  status = curandGenerate(randGenerator, outputPtr, num);
+
 #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;

diff --git a/tests/unit_tests/synthetic/libraries/curand2rocrand.cu b/tests/unit_tests/synthetic/libraries/curand2rocrand.cu
@@ -10,8 +10,11 @@
 int main() {
   printf("21.1. cuRAND API to rocRAND API synthetic test\n");
 
+  unsigned int *outputPtr = nullptr;
+  size_t num = 0;
+
   // CHECK: rocrand_status randStatus;
-  // CHECK-NEXT: rocrand_status randStatus_t;
+  // CHECK-NEXT: rocrand_status status;
   // CHECK-NEXT: rocrand_status STATUS_SUCCESS = ROCRAND_STATUS_SUCCESS;
   // CHECK-NEXT: rocrand_status STATUS_VERSION_MISMATCH = ROCRAND_STATUS_VERSION_MISMATCH;
   // CHECK-NEXT: rocrand_status STATUS_NOT_INITIALIZED = ROCRAND_STATUS_NOT_CREATED;
@@ -23,7 +26,7 @@ int main() {
   // CHECK-NEXT: rocrand_status STATUS_LAUNCH_FAILURE = ROCRAND_STATUS_LAUNCH_FAILURE;
   // CHECK-NEXT: rocrand_status STATUS_INTERNAL_ERROR = ROCRAND_STATUS_INTERNAL_ERROR;
   curandStatus randStatus;
-  curandStatus_t randStatus_t;
+  curandStatus_t status;
   curandStatus_t STATUS_SUCCESS = CURAND_STATUS_SUCCESS;
   curandStatus_t STATUS_VERSION_MISMATCH = CURAND_STATUS_VERSION_MISMATCH;
   curandStatus_t STATUS_NOT_INITIALIZED = CURAND_STATUS_NOT_INITIALIZED;
@@ -92,6 +95,26 @@ int main() {
   curandGenerator_st *randGenerator_st = nullptr;
   curandGenerator_t randGenerator;
 
+  // CUDA: curandStatus_t CURANDAPI curandCreateGenerator(curandGenerator_t *generator, curandRngType_t rng_type);
+  // ROC: rocrand_status ROCRANDAPI rocrand_create_generator(rocrand_generator * generator, rocrand_rng_type rng_type);
+  // CHECK: status = rocrand_create_generator(&randGenerator, randRngType_t);
+  status = curandCreateGenerator(&randGenerator, randRngType_t);
+
+  // CUDA: curandStatus_t CURANDAPI curandDestroyGenerator(curandGenerator_t generator);
+  // ROC: rocrand_status ROCRANDAPI rocrand_destroy_generator(rocrand_generator generator);
+  // CHECK: status = rocrand_destroy_generator(randGenerator);
+  status = curandDestroyGenerator(randGenerator);
+
+  // CUDA: curandStatus_t CURANDAPI curandCreateGeneratorHost(curandGenerator_t *generator, curandRngType_t rng_type);
+  // ROC: rocrand_status ROCRANDAPI rocrand_create_generator_host_blocking(rocrand_generator* generator, rocrand_rng_type rng_type);
+  // CHECK: status = rocrand_create_generator_host_blocking(&randGenerator, randRngType_t);
+  status = curandCreateGeneratorHost(&randGenerator, randRngType_t);
+
+  // CUDA: curandStatus_t CURANDAPI curandGenerate(curandGenerator_t generator, unsigned int *outputPtr, size_t num);
+  // ROC: rocrand_status ROCRANDAPI rocrand_generate(rocrand_generator generator, unsigned int * output_data, size_t n);
+  // CHECK: status = rocrand_generate(randGenerator, outputPtr, num);
+  status = curandGenerate(randGenerator, outputPtr, num);
+
 #if CUDA_VERSION >= 11000 && CURAND_VERSION >= 10200
   // CHECK: rocrand_ordering RAND_ORDERING_PSEUDO_LEGACY = ROCRAND_ORDERING_PSEUDO_LEGACY;
   curandOrdering_t RAND_ORDERING_PSEUDO_LEGACY = CURAND_ORDERING_PSEUDO_LEGACY;