[HIPIFY][6.0.0][hipRAND][doc] Support for ROCm HIP 6.0.0 - Step 32 - …

…hipRAND + Updated the `benchmark_curand_kernel.cpp` test, the regenerated hipify-perl and RAND docs [TODO] + Add synthetic tests on RAND
emankov · Oct 30, 2023 · 788f815 · 788f815
1 parent 328466d
commit 788f815
Show file tree

Hide file tree

Showing 5 changed files with 40 additions and 32 deletions.
diff --git a/bin/hipify-perl b/bin/hipify-perl
@@ -3504,6 +3504,10 @@ sub simpleSubstitutions {
     subst("curandGenerateSeeds", "hiprandGenerateSeeds", "library");
     subst("curandGenerateUniform", "hiprandGenerateUniform", "library");
     subst("curandGenerateUniformDouble", "hiprandGenerateUniformDouble", "library");
+    subst("curandGetDirectionVectors32", "hiprandGetDirectionVectors32", "library");
+    subst("curandGetDirectionVectors64", "hiprandGetDirectionVectors64", "library");
+    subst("curandGetScrambleConstants32", "hiprandGetScrambleConstants32", "library");
+    subst("curandGetScrambleConstants64", "hiprandGetScrambleConstants64", "library");
     subst("curandGetVersion", "hiprandGetVersion", "library");
     subst("curandMakeMTGP32Constants", "hiprandMakeMTGP32Constants", "library");
     subst("curandMakeMTGP32KernelState", "hiprandMakeMTGP32KernelState", "library");
@@ -4543,6 +4547,8 @@ sub simpleSubstitutions {
     subst("cufftType_t", "hipfftType_t", "type");
     subst("cufftXtCallbackType", "hipfftXtCallbackType", "type");
     subst("cufftXtCallbackType_t", "hipfftXtCallbackType_t", "type");
+    subst("curandDirectionVectorSet", "hiprandDirectionVectorSet_t", "type");
+    subst("curandDirectionVectorSet_t", "hiprandDirectionVectorSet_t", "type");
     subst("curandDirectionVectors32_t", "hiprandDirectionVectors32_t", "type");
     subst("curandDiscreteDistribution_st", "hiprandDiscreteDistribution_st", "type");
     subst("curandDiscreteDistribution_t", "hiprandDiscreteDistribution_t", "type");
@@ -4883,6 +4889,8 @@ sub simpleSubstitutions {
     subst("CUFFT_UNALIGNED_DATA", "HIPFFT_UNALIGNED_DATA", "numeric_literal");
     subst("CUFFT_Z2D", "HIPFFT_Z2D", "numeric_literal");
     subst("CUFFT_Z2Z", "HIPFFT_Z2Z", "numeric_literal");
+    subst("CURAND_DIRECTION_VECTORS_32_JOEKUO6", "HIPRAND_DIRECTION_VECTORS_32_JOEKUO6", "numeric_literal");
+    subst("CURAND_DIRECTION_VECTORS_64_JOEKUO6", "HIPRAND_DIRECTION_VECTORS_64_JOEKUO6", "numeric_literal");
     subst("CURAND_RNG_PSEUDO_DEFAULT", "HIPRAND_RNG_PSEUDO_DEFAULT", "numeric_literal");
     subst("CURAND_RNG_PSEUDO_MRG32K3A", "HIPRAND_RNG_PSEUDO_MRG32K3A", "numeric_literal");
     subst("CURAND_RNG_PSEUDO_MT19937", "HIPRAND_RNG_PSEUDO_MT19937", "numeric_literal");
@@ -4895,6 +4903,8 @@ sub simpleSubstitutions {
     subst("CURAND_RNG_QUASI_SOBOL32", "HIPRAND_RNG_QUASI_SOBOL32", "numeric_literal");
     subst("CURAND_RNG_QUASI_SOBOL64", "HIPRAND_RNG_QUASI_SOBOL64", "numeric_literal");
     subst("CURAND_RNG_TEST", "HIPRAND_RNG_TEST", "numeric_literal");
+    subst("CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", "HIPRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", "numeric_literal");
+    subst("CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", "HIPRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", "numeric_literal");
     subst("CURAND_STATUS_ALLOCATION_FAILED", "HIPRAND_STATUS_ALLOCATION_FAILED", "numeric_literal");
     subst("CURAND_STATUS_ARCH_MISMATCH", "HIPRAND_STATUS_ARCH_MISMATCH", "numeric_literal");
     subst("CURAND_STATUS_DOUBLE_PRECISION_REQUIRED", "HIPRAND_STATUS_DOUBLE_PRECISION_REQUIRED", "numeric_literal");
@@ -6970,11 +6980,7 @@ sub warnUnsupportedFunctions {
         "curandHistogramM2V_st",
         "curandHistogramM2K_t",
         "curandHistogramM2K_st",
-        "curandGetScrambleConstants64",
-        "curandGetScrambleConstants32",
         "curandGetProperty",
-        "curandGetDirectionVectors64",
-        "curandGetDirectionVectors32",
         "curandGenerateLongLong",
         "curandDistribution_t",
         "curandDistribution_st",
@@ -6983,8 +6989,6 @@ sub warnUnsupportedFunctions {
         "curandDistributionM2Shift_t",
         "curandDistributionM2Shift_st",
         "curandDirectionVectors64_t",
-        "curandDirectionVectorSet_t",
-        "curandDirectionVectorSet",
         "cufftXtWorkAreaPolicy_t",
         "cufftXtWorkAreaPolicy",
         "cufftXtSubFormat_t",
@@ -8563,8 +8567,6 @@ sub warnUnsupportedFunctions {
         "CUSPARSE_ALG_MERGE_PATH",
         "CUSPARSE_ALG1",
         "CUSPARSE_ALG0",
-        "CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6",
-        "CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6",
         "CURAND_REJECTION",
         "CURAND_POISSON",
         "CURAND_ORDERING_QUASI_DEFAULT",
@@ -8580,8 +8582,6 @@ sub warnUnsupportedFunctions {
         "CURAND_HITR",
         "CURAND_FAST_REJECTION",
         "CURAND_DISCRETE_GAUSS",
-        "CURAND_DIRECTION_VECTORS_64_JOEKUO6",
-        "CURAND_DIRECTION_VECTORS_32_JOEKUO6",
         "CURAND_DEVICE_API",
         "CURAND_DEFINITION",
         "CURAND_CHOOSE_BEST",

diff --git a/docs/tables/CURAND_API_supported_by_HIP.md b/docs/tables/CURAND_API_supported_by_HIP.md
@@ -9,8 +9,8 @@
 |`CURAND_CHOOSE_BEST`| | | | | | | | | | |
 |`CURAND_DEFINITION`| | | | | | | | | | |
 |`CURAND_DEVICE_API`| | | | | | | | | | |
-|`CURAND_DIRECTION_VECTORS_32_JOEKUO6`| | | | | | | | | | |
-|`CURAND_DIRECTION_VECTORS_64_JOEKUO6`| | | | | | | | | | |
+|`CURAND_DIRECTION_VECTORS_32_JOEKUO6`| | | | |`HIPRAND_DIRECTION_VECTORS_32_JOEKUO6`|6.0.0| | | |6.0.0|
+|`CURAND_DIRECTION_VECTORS_64_JOEKUO6`| | | | |`HIPRAND_DIRECTION_VECTORS_64_JOEKUO6`|6.0.0| | | |6.0.0|
 |`CURAND_DISCRETE_GAUSS`| | | | | | | | | | |
 |`CURAND_FAST_REJECTION`| | | | | | | | | | |
 |`CURAND_HITR`| | | | | | | | | | |
@@ -38,8 +38,8 @@
 |`CURAND_RNG_QUASI_SOBOL32`| | | | |`HIPRAND_RNG_QUASI_SOBOL32`|1.5.0| | | | |
 |`CURAND_RNG_QUASI_SOBOL64`| | | | |`HIPRAND_RNG_QUASI_SOBOL64`|1.5.0| | | | |
 |`CURAND_RNG_TEST`| | | | |`HIPRAND_RNG_TEST`|1.5.0| | | | |
-|`CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`| | | | | | | | | | |
-|`CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`| | | | | | | | | | |
+|`CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`| | | | |`HIPRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`|6.0.0| | | |6.0.0|
+|`CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`| | | | |`HIPRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`|6.0.0| | | |6.0.0|
 |`CURAND_STATUS_ALLOCATION_FAILED`| | | | |`HIPRAND_STATUS_ALLOCATION_FAILED`|1.5.0| | | | |
 |`CURAND_STATUS_ARCH_MISMATCH`| | | | |`HIPRAND_STATUS_ARCH_MISMATCH`|1.5.0| | | | |
 |`CURAND_STATUS_DOUBLE_PRECISION_REQUIRED`| | | | |`HIPRAND_STATUS_DOUBLE_PRECISION_REQUIRED`|1.5.0| | | | |
@@ -53,8 +53,8 @@
 |`CURAND_STATUS_SUCCESS`| | | | |`HIPRAND_STATUS_SUCCESS`|1.5.0| | | | |
 |`CURAND_STATUS_TYPE_ERROR`| | | | |`HIPRAND_STATUS_TYPE_ERROR`|1.5.0| | | | |
 |`CURAND_STATUS_VERSION_MISMATCH`| | | | |`HIPRAND_STATUS_VERSION_MISMATCH`|1.5.0| | | | |
-|`curandDirectionVectorSet`| | | | | | | | | | |
-|`curandDirectionVectorSet_t`| | | | | | | | | | |
+|`curandDirectionVectorSet`| | | | |`hiprandDirectionVectorSet_t`|6.0.0| | | |6.0.0|
+|`curandDirectionVectorSet_t`| | | | |`hiprandDirectionVectorSet_t`|6.0.0| | | |6.0.0|
 |`curandDirectionVectors32_t`| | | | |`hiprandDirectionVectors32_t`|1.5.0| | | | |
 |`curandDirectionVectors64_t`| | | | | | | | | | |
 |`curandDiscreteDistribution_st`| | | | |`hiprandDiscreteDistribution_st`|1.5.0| | | | |
@@ -119,11 +119,11 @@
 |`curandGenerateSeeds`| | | | |`hiprandGenerateSeeds`|1.5.0| | | | |
 |`curandGenerateUniform`| | | | |`hiprandGenerateUniform`|1.5.0| | | | |
 |`curandGenerateUniformDouble`| | | | |`hiprandGenerateUniformDouble`|1.5.0| | | | |
-|`curandGetDirectionVectors32`| | | | | | | | | | |
-|`curandGetDirectionVectors64`| | | | | | | | | | |
+|`curandGetDirectionVectors32`| | | | |`hiprandGetDirectionVectors32`|6.0.0| | | |6.0.0|
+|`curandGetDirectionVectors64`| | | | |`hiprandGetDirectionVectors64`|6.0.0| | | |6.0.0|
 |`curandGetProperty`|8.0| | | | | | | | | |
-|`curandGetScrambleConstants32`| | | | | | | | | | |
-|`curandGetScrambleConstants64`| | | | | | | | | | |
+|`curandGetScrambleConstants32`| | | | |`hiprandGetScrambleConstants32`|6.0.0| | | |6.0.0|
+|`curandGetScrambleConstants64`| | | | |`hiprandGetScrambleConstants64`|6.0.0| | | |6.0.0|
 |`curandGetVersion`| | | | |`hiprandGetVersion`|1.5.0| | | | |
 |`curandMakeMTGP32Constants`| | | | |`hiprandMakeMTGP32Constants`|1.5.0| | | | |
 |`curandMakeMTGP32KernelState`| | | | |`hiprandMakeMTGP32KernelState`|1.5.0| | | | |

diff --git a/src/CUDA2HIP_RAND_API_functions.cpp b/src/CUDA2HIP_RAND_API_functions.cpp
@@ -40,11 +40,11 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RAND_FUNCTION_MAP {
   {"curandGenerateSeeds",                           {"hiprandGenerateSeeds",                           "", CONV_LIB_FUNC, API_RAND, 2}},
   {"curandGenerateUniform",                         {"hiprandGenerateUniform",                         "", CONV_LIB_FUNC, API_RAND, 2}},
   {"curandGenerateUniformDouble",                   {"hiprandGenerateUniformDouble",                   "", CONV_LIB_FUNC, API_RAND, 2}},
-  {"curandGetDirectionVectors32",                   {"hiprandGetDirectionVectors32",                   "", CONV_LIB_FUNC, API_RAND, 2, HIP_UNSUPPORTED}},
-  {"curandGetDirectionVectors64",                   {"hiprandGetDirectionVectors64",                   "", CONV_LIB_FUNC, API_RAND, 2, HIP_UNSUPPORTED}},
+  {"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, HIP_UNSUPPORTED}},
-  {"curandGetScrambleConstants64",                  {"hiprandGetScrambleConstants64",                  "", 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}},
   {"curandSetGeneratorOffset",                      {"hiprandSetGeneratorOffset",                      "", CONV_LIB_FUNC, API_RAND, 2}},
   {"curandSetGeneratorOrdering",                    {"hiprandSetGeneratorOrdering",                    "", CONV_LIB_FUNC, API_RAND, 2, HIP_UNSUPPORTED}},
@@ -136,6 +136,10 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_RAND_FUNCTION_VER_MAP {
   {"hiprand_discrete4",                              {HIP_1050, HIP_0,    HIP_0   }},
   {"hiprand_poisson",                                {HIP_1050, HIP_0,    HIP_0   }},
   {"hiprand_poisson4",                               {HIP_1050, HIP_0,    HIP_0   }},
+  {"hiprandGetDirectionVectors32",                   {HIP_6000, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"hiprandGetDirectionVectors64",                   {HIP_6000, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"hiprandGetScrambleConstants32",                  {HIP_6000, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"hiprandGetScrambleConstants64",                  {HIP_6000, HIP_0,    HIP_0,  HIP_LATEST}},
 };
 
 const std::map<unsigned int, llvm::StringRef> CUDA_RAND_API_SECTION_MAP {

diff --git a/src/CUDA2HIP_RAND_API_types.cpp b/src/CUDA2HIP_RAND_API_types.cpp
@@ -31,8 +31,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RAND_TYPE_NAME_MAP {
   {"curandRngType_t",               {"hiprandRngType_t",               "", CONV_TYPE, API_RAND, 1}},
   {"curandGenerator_st",            {"hiprandGenerator_st",            "", CONV_TYPE, API_RAND, 1}},
   {"curandGenerator_t",             {"hiprandGenerator_t",             "", CONV_TYPE, API_RAND, 1}},
-  {"curandDirectionVectorSet",      {"hiprandDirectionVectorSet_t",    "", CONV_TYPE, API_RAND, 1, HIP_UNSUPPORTED}},
-  {"curandDirectionVectorSet_t",    {"hiprandDirectionVectorSet_t",    "", CONV_TYPE, API_RAND, 1, HIP_UNSUPPORTED}},
+  {"curandDirectionVectorSet",      {"hiprandDirectionVectorSet_t",    "", CONV_TYPE, API_RAND, 1}},
+  {"curandDirectionVectorSet_t",    {"hiprandDirectionVectorSet_t",    "", CONV_TYPE, API_RAND, 1}},
   {"curandOrdering",                {"hiprandOrdering_t",              "", CONV_TYPE, API_RAND, 1, HIP_UNSUPPORTED}},
   {"curandOrdering_t",              {"hiprandOrdering_t",              "", CONV_TYPE, API_RAND, 1, HIP_UNSUPPORTED}},
   {"curandDistribution_st",         {"hiprandDistribution_st",         "", CONV_TYPE, API_RAND, 1, HIP_UNSUPPORTED}},
@@ -112,10 +112,10 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RAND_TYPE_NAME_MAP {
   {"CURAND_ORDERING_QUASI_DEFAULT",                 {"HIPRAND_ORDERING_QUASI_DEFAULT",                 "", CONV_NUMERIC_LITERAL, API_RAND, 1, HIP_UNSUPPORTED}},
 
   // RAND choice of direction vector set (enum curandDirectionVectorSet)
-  {"CURAND_DIRECTION_VECTORS_32_JOEKUO6",           {"HIPRAND_DIRECTION_VECTORS_32_JOEKUO6",           "", CONV_NUMERIC_LITERAL, API_RAND, 1, HIP_UNSUPPORTED}},
-  {"CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", {"HIPRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", "", CONV_NUMERIC_LITERAL, API_RAND, 1, HIP_UNSUPPORTED}},
-  {"CURAND_DIRECTION_VECTORS_64_JOEKUO6",           {"HIPRAND_DIRECTION_VECTORS_64_JOEKUO6",           "", CONV_NUMERIC_LITERAL, API_RAND, 1, HIP_UNSUPPORTED}},
-  {"CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", {"HIPRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", "", CONV_NUMERIC_LITERAL, API_RAND, 1, HIP_UNSUPPORTED}},
+  {"CURAND_DIRECTION_VECTORS_32_JOEKUO6",           {"HIPRAND_DIRECTION_VECTORS_32_JOEKUO6",           "", CONV_NUMERIC_LITERAL, API_RAND, 1}},
+  {"CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", {"HIPRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", "", CONV_NUMERIC_LITERAL, API_RAND, 1}},
+  {"CURAND_DIRECTION_VECTORS_64_JOEKUO6",           {"HIPRAND_DIRECTION_VECTORS_64_JOEKUO6",           "", CONV_NUMERIC_LITERAL, API_RAND, 1}},
+  {"CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", {"HIPRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", "", CONV_NUMERIC_LITERAL, API_RAND, 1}},
 
   // RAND method (enum curandMethod)
   {"CURAND_CHOOSE_BEST",                            {"HIPRAND_CHOOSE_BEST",                            "", CONV_NUMERIC_LITERAL, API_RAND, 1, HIP_UNSUPPORTED}},
@@ -184,4 +184,9 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_RAND_TYPE_NAME_VER_MAP {
   {"HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL32",            {HIP_1050, HIP_0,    HIP_0   }},
   {"HIPRAND_RNG_QUASI_SOBOL64",                      {HIP_1050, HIP_0,    HIP_0   }},
   {"HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64",            {HIP_1050, HIP_0,    HIP_0   }},
+  {"hiprandDirectionVectorSet_t",                    {HIP_6000, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"HIPRAND_DIRECTION_VECTORS_32_JOEKUO6",           {HIP_6000, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"HIPRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", {HIP_6000, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"HIPRAND_DIRECTION_VECTORS_64_JOEKUO6",           {HIP_6000, HIP_0,    HIP_0,  HIP_LATEST}},
+  {"HIPRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", {HIP_6000, HIP_0,    HIP_0,  HIP_LATEST}},
 };
diff --git a/tests/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp b/tests/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp
@@ -297,8 +297,7 @@ struct runner<curandStateSobol32_t>
         CUDA_CALL(cudaMalloc((void **)&directions, size));
         // CHECK: hiprandDirectionVectors32_t * h_directions;
         curandDirectionVectors32_t * h_directions;
-        // hiprandGetDirectionVectors32 and HIPRAND_DIRECTION_VECTORS_32_JOEKUO6 (of hiprandDirectionVectorSet_t) are yet unsupported by HIP
-        // CHECK-NOT: CURAND_CALL(hiprandGetDirectionVectors32(&h_directions, HIPRAND_DIRECTION_VECTORS_32_JOEKUO6));
+        // CHECK: CURAND_CALL(hiprandGetDirectionVectors32(&h_directions, HIPRAND_DIRECTION_VECTORS_32_JOEKUO6));
         CURAND_CALL(curandGetDirectionVectors32(&h_directions, CURAND_DIRECTION_VECTORS_32_JOEKUO6));
         // CHECK: CUDA_CALL(hipMemcpy(directions, h_directions, size, hipMemcpyHostToDevice));
         CUDA_CALL(cudaMemcpy(directions, h_directions, size, cudaMemcpyHostToDevice));