From 187080bb0bc799f1c4e4a579f171c87b3bde0012 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 16 Sep 2024 16:49:09 +0100 Subject: [PATCH] [HIPIFY][rocRAND][feature] Support for `cuRAND -> rocRAND` hipification - Step 7 - Functions + Updated synthetic tests for `rocRAND` and `hipRAND`, the regenerated `hipify-perl`, and `RAND` `CUDA2HIP` documentation --- bin/hipify-perl | 4 +++ .../CURAND_API_supported_by_HIP_and_ROC.md | 8 +++--- docs/tables/CURAND_API_supported_by_ROC.md | 8 +++--- src/CUDA2HIP_RAND_API_functions.cpp | 13 ++++++--- .../synthetic/libraries/curand2hiprand.cu | 23 ++++++++++++++++ .../synthetic/libraries/curand2rocrand.cu | 27 +++++++++++++++++-- 6 files changed, 69 insertions(+), 14 deletions(-) diff --git a/bin/hipify-perl b/bin/hipify-perl index ab8f80cf..304553bf 100755 --- 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 index a3290081..b88ff118 100644 --- 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 index 2ad59fb2..c673fc40 100644 --- 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 index 89c5f04d..8f48ded9 100644 --- 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 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 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 CUDA_RAND_API_SECTION_MAP { diff --git a/tests/unit_tests/synthetic/libraries/curand2hiprand.cu b/tests/unit_tests/synthetic/libraries/curand2hiprand.cu index 934a962c..cec61e9a 100644 --- 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 index 1a851135..f54e5e4b 100644 --- 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;