[HIPIFY][rocRAND][feature] Support for cuRAND -> rocRAND hipificati…

…on - Step 6 - Data Types + Updated synthetic tests for `rocRAND` and `hipRAND`, the regenerated `hipify-perl`, and `RAND` `CUDA2HIP` documentation
emankov · Sep 11, 2024 · 4d0d40d · 4d0d40d
1 parent c678c5a
commit 4d0d40d
Show file tree

Hide file tree

Showing 6 changed files with 77 additions and 32 deletions.
diff --git a/bin/hipify-perl b/bin/hipify-perl
@@ -3207,6 +3207,10 @@ sub rocSubstitutions {
     subst("cudnnStatus_t", "miopenStatus_t", "type");
     subst("cudnnTensorDescriptor_t", "miopenTensorDescriptor_t", "type");
     subst("cudnnTensorFormat_t", "miopenTensorLayout_t", "type");
+    subst("curandDirectionVectorSet", "rocrand_direction_vector_set", "type");
+    subst("curandDirectionVectorSet_t", "rocrand_direction_vector_set", "type");
+    subst("curandGenerator_st", "rocrand_generator_base_type", "type");
+    subst("curandGenerator_t", "rocrand_generator", "type");
     subst("curandOrdering", "rocrand_ordering", "type");
     subst("curandOrdering_t", "rocrand_ordering", "type");
     subst("curandRngType", "rocrand_rng_type", "type");
@@ -3653,6 +3657,8 @@ sub rocSubstitutions {
     subst("CUDNN_TYPE_TENSOR_REORDERING_MODE", "MIOPEN_TYPE_TENSOR_REORDERING_MODE", "numeric_literal");
     subst("CUDNN_TYPE_VOID_PTR", "MIOPEN_TYPE_VOID_PTR", "numeric_literal");
     subst("CUDNN_UNIDIRECTIONAL", "miopenRNNunidirection", "numeric_literal");
+    subst("CURAND_DIRECTION_VECTORS_32_JOEKUO6", "ROCRAND_DIRECTION_VECTORS_32_JOEKUO6", "numeric_literal");
+    subst("CURAND_DIRECTION_VECTORS_64_JOEKUO6", "ROCRAND_DIRECTION_VECTORS_64_JOEKUO6", "numeric_literal");
     subst("CURAND_ORDERING_PSEUDO_BEST", "ROCRAND_ORDERING_PSEUDO_BEST", "numeric_literal");
     subst("CURAND_ORDERING_PSEUDO_DEFAULT", "ROCRAND_ORDERING_PSEUDO_DEFAULT", "numeric_literal");
     subst("CURAND_ORDERING_PSEUDO_DYNAMIC", "ROCRAND_ORDERING_PSEUDO_DYNAMIC", "numeric_literal");
@@ -3670,6 +3676,8 @@ sub rocSubstitutions {
     subst("CURAND_RNG_QUASI_SCRAMBLED_SOBOL64", "ROCRAND_RNG_QUASI_SCRAMBLED_SOBOL64", "numeric_literal");
     subst("CURAND_RNG_QUASI_SOBOL32", "ROCRAND_RNG_QUASI_SOBOL32", "numeric_literal");
     subst("CURAND_RNG_QUASI_SOBOL64", "ROCRAND_RNG_QUASI_SOBOL64", "numeric_literal");
+    subst("CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", "ROCRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", "numeric_literal");
+    subst("CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", "ROCRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", "numeric_literal");
     subst("CURAND_STATUS_ALLOCATION_FAILED", "ROCRAND_STATUS_ALLOCATION_FAILED", "numeric_literal");
     subst("CURAND_STATUS_DOUBLE_PRECISION_REQUIRED", "ROCRAND_STATUS_DOUBLE_PRECISION_REQUIRED", "numeric_literal");
     subst("CURAND_STATUS_INTERNAL_ERROR", "ROCRAND_STATUS_INTERNAL_ERROR", "numeric_literal");

diff --git a/docs/tables/CURAND_API_supported_by_HIP_and_ROC.md b/docs/tables/CURAND_API_supported_by_HIP_and_ROC.md
@@ -9,8 +9,8 @@
 |`CURAND_CHOOSE_BEST`| | | | | | | | | | | | | | | | |
 |`CURAND_DEFINITION`| | | | | | | | | | | | | | | | |
 |`CURAND_DEVICE_API`| | | | | | | | | | | | | | | | |
-|`CURAND_DIRECTION_VECTORS_32_JOEKUO6`| | | | |`HIPRAND_DIRECTION_VECTORS_32_JOEKUO6`|6.0.0| | | | | | | | | | |
-|`CURAND_DIRECTION_VECTORS_64_JOEKUO6`| | | | |`HIPRAND_DIRECTION_VECTORS_64_JOEKUO6`|6.0.0| | | | | | | | | | |
+|`CURAND_DIRECTION_VECTORS_32_JOEKUO6`| | | | |`HIPRAND_DIRECTION_VECTORS_32_JOEKUO6`|6.0.0| | | | |`ROCRAND_DIRECTION_VECTORS_32_JOEKUO6`|6.0.0| | | | |
+|`CURAND_DIRECTION_VECTORS_64_JOEKUO6`| | | | |`HIPRAND_DIRECTION_VECTORS_64_JOEKUO6`|6.0.0| | | | |`ROCRAND_DIRECTION_VECTORS_64_JOEKUO6`|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| | | | |`ROCRAND_RNG_QUASI_SOBOL32`|1.5.0| | | | |
 |`CURAND_RNG_QUASI_SOBOL64`| | | | |`HIPRAND_RNG_QUASI_SOBOL64`|1.5.0| | | | |`ROCRAND_RNG_QUASI_SOBOL64`|4.5.0| | | | |
 |`CURAND_RNG_TEST`| | | | |`HIPRAND_RNG_TEST`|1.5.0| | | | | | | | | | |
-|`CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`| | | | |`HIPRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`|6.0.0| | | | | | | | | | |
-|`CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`| | | | |`HIPRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`|6.0.0| | | | | | | | | | |
+|`CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`| | | | |`HIPRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`|6.0.0| | | | |`ROCRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`|6.0.0| | | | |
+|`CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`| | | | |`HIPRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`|6.0.0| | | | |`ROCRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`|6.0.0| | | | |
 |`CURAND_STATUS_ALLOCATION_FAILED`| | | | |`HIPRAND_STATUS_ALLOCATION_FAILED`|1.5.0| | | | |`ROCRAND_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| | | | |`ROCRAND_STATUS_DOUBLE_PRECISION_REQUIRED`|1.5.0| | | | |
@@ -53,8 +53,8 @@
 |`CURAND_STATUS_SUCCESS`| | | | |`HIPRAND_STATUS_SUCCESS`|1.5.0| | | | |`ROCRAND_STATUS_SUCCESS`|1.5.0| | | | |
 |`CURAND_STATUS_TYPE_ERROR`| | | | |`HIPRAND_STATUS_TYPE_ERROR`|1.5.0| | | | |`ROCRAND_STATUS_TYPE_ERROR`|1.5.0| | | | |
 |`CURAND_STATUS_VERSION_MISMATCH`| | | | |`HIPRAND_STATUS_VERSION_MISMATCH`|1.5.0| | | | |`ROCRAND_STATUS_VERSION_MISMATCH`|1.5.0| | | | |
-|`curandDirectionVectorSet`| | | | |`hiprandDirectionVectorSet_t`|6.0.0| | | | | | | | | | |
-|`curandDirectionVectorSet_t`| | | | |`hiprandDirectionVectorSet_t`|6.0.0| | | | | | | | | | |
+|`curandDirectionVectorSet`| | | | |`hiprandDirectionVectorSet_t`|6.0.0| | | | |`rocrand_direction_vector_set`|6.0.0| | | | |
+|`curandDirectionVectorSet_t`| | | | |`hiprandDirectionVectorSet_t`|6.0.0| | | | |`rocrand_direction_vector_set`|6.0.0| | | | |
 |`curandDirectionVectors32_t`| | | | |`hiprandDirectionVectors32_t`|1.5.0| | | | | | | | | | |
 |`curandDirectionVectors64_t`| | | | |`hiprandDirectionVectors64_t`|6.0.0| | | | | | | | | | |
 |`curandDiscreteDistribution_st`| | | | |`hiprandDiscreteDistribution_st`|1.5.0| | | | | | | | | | |
@@ -65,8 +65,8 @@
 |`curandDistributionShift_t`| | | | | | | | | | | | | | | | |
 |`curandDistribution_st`| | | | | | | | | | | | | | | | |
 |`curandDistribution_t`| | | | | | | | | | | | | | | | |
-|`curandGenerator_st`| | | | |`hiprandGenerator_st`|1.5.0| | | | | | | | | | |
-|`curandGenerator_t`| | | | |`hiprandGenerator_t`|1.5.0| | | | | | | | | | |
+|`curandGenerator_st`| | | | |`hiprandGenerator_st`|1.5.0| | | | |`rocrand_generator_base_type`|1.5.0| | | | |
+|`curandGenerator_t`| | | | |`hiprandGenerator_t`|1.5.0| | | | |`rocrand_generator`|1.5.0| | | | |
 |`curandHistogramM2K_st`| | | | | | | | | | | | | | | | |
 |`curandHistogramM2K_t`| | | | | | | | | | | | | | | | |
 |`curandHistogramM2V_st`| | | | | | | | | | | | | | | | |

diff --git a/docs/tables/CURAND_API_supported_by_ROC.md b/docs/tables/CURAND_API_supported_by_ROC.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`| | | | |`ROCRAND_DIRECTION_VECTORS_32_JOEKUO6`|6.0.0| | | | |
+|`CURAND_DIRECTION_VECTORS_64_JOEKUO6`| | | | |`ROCRAND_DIRECTION_VECTORS_64_JOEKUO6`|6.0.0| | | | |
 |`CURAND_DISCRETE_GAUSS`| | | | | | | | | | |
 |`CURAND_FAST_REJECTION`| | | | | | | | | | |
 |`CURAND_HITR`| | | | | | | | | | |
@@ -38,8 +38,8 @@
 |`CURAND_RNG_QUASI_SOBOL32`| | | | |`ROCRAND_RNG_QUASI_SOBOL32`|1.5.0| | | | |
 |`CURAND_RNG_QUASI_SOBOL64`| | | | |`ROCRAND_RNG_QUASI_SOBOL64`|4.5.0| | | | |
 |`CURAND_RNG_TEST`| | | | | | | | | | |
-|`CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`| | | | | | | | | | |
-|`CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`| | | | | | | | | | |
+|`CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`| | | | |`ROCRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6`|6.0.0| | | | |
+|`CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`| | | | |`ROCRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6`|6.0.0| | | | |
 |`CURAND_STATUS_ALLOCATION_FAILED`| | | | |`ROCRAND_STATUS_ALLOCATION_FAILED`|1.5.0| | | | |
 |`CURAND_STATUS_ARCH_MISMATCH`| | | | | | | | | | |
 |`CURAND_STATUS_DOUBLE_PRECISION_REQUIRED`| | | | |`ROCRAND_STATUS_DOUBLE_PRECISION_REQUIRED`|1.5.0| | | | |
@@ -53,8 +53,8 @@
 |`CURAND_STATUS_SUCCESS`| | | | |`ROCRAND_STATUS_SUCCESS`|1.5.0| | | | |
 |`CURAND_STATUS_TYPE_ERROR`| | | | |`ROCRAND_STATUS_TYPE_ERROR`|1.5.0| | | | |
 |`CURAND_STATUS_VERSION_MISMATCH`| | | | |`ROCRAND_STATUS_VERSION_MISMATCH`|1.5.0| | | | |
-|`curandDirectionVectorSet`| | | | | | | | | | |
-|`curandDirectionVectorSet_t`| | | | | | | | | | |
+|`curandDirectionVectorSet`| | | | |`rocrand_direction_vector_set`|6.0.0| | | | |
+|`curandDirectionVectorSet_t`| | | | |`rocrand_direction_vector_set`|6.0.0| | | | |
 |`curandDirectionVectors32_t`| | | | | | | | | | |
 |`curandDirectionVectors64_t`| | | | | | | | | | |
 |`curandDiscreteDistribution_st`| | | | | | | | | | |
@@ -65,8 +65,8 @@
 |`curandDistributionShift_t`| | | | | | | | | | |
 |`curandDistribution_st`| | | | | | | | | | |
 |`curandDistribution_t`| | | | | | | | | | |
-|`curandGenerator_st`| | | | | | | | | | |
-|`curandGenerator_t`| | | | | | | | | | |
+|`curandGenerator_st`| | | | |`rocrand_generator_base_type`|1.5.0| | | | |
+|`curandGenerator_t`| | | | |`rocrand_generator`|1.5.0| | | | |
 |`curandHistogramM2K_st`| | | | | | | | | | |
 |`curandHistogramM2K_t`| | | | | | | | | | |
 |`curandHistogramM2V_st`| | | | | | | | | | |

diff --git a/src/CUDA2HIP_RAND_API_types.cpp b/src/CUDA2HIP_RAND_API_types.cpp
@@ -68,16 +68,26 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RAND_TYPE_NAME_MAP {
   {"CURAND_ORDERING_PSEUDO_DYNAMIC",                {"HIPRAND_ORDERING_PSEUDO_DYNAMIC",                "ROCRAND_ORDERING_PSEUDO_DYNAMIC",                                CONV_NUMERIC_LITERAL, API_RAND, 1, HIP_EXPERIMENTAL}},
   {"CURAND_ORDERING_QUASI_DEFAULT",                 {"HIPRAND_ORDERING_QUASI_DEFAULT",                 "ROCRAND_ORDERING_QUASI_DEFAULT",                                 CONV_NUMERIC_LITERAL, API_RAND, 1, HIP_EXPERIMENTAL}},
 
-  {"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}},
-  {"curandDirectionVectorSet_t",                    {"hiprandDirectionVectorSet_t",                    "", CONV_TYPE, API_RAND, 1}},
-  {"curandDistribution_st",                         {"hiprandDistribution_st",                         "", CONV_TYPE, API_RAND, 1, HIP_UNSUPPORTED}},
-  {"curandHistogramM2V_st",                         {"hiprandHistogramM2V_st",                         "", CONV_TYPE, API_RAND, 1, HIP_UNSUPPORTED}},
-  {"curandDistribution_t",                          {"hiprandDistribution_t",                          "", CONV_TYPE, API_RAND, 1, HIP_UNSUPPORTED}},
-  {"curandHistogramM2V_t",                          {"hiprandDistribution_t",                          "", CONV_TYPE, API_RAND, 1, HIP_UNSUPPORTED}},
-  {"curandDistributionShift_st",                    {"hiprandDistributionShift_st",                    "", CONV_TYPE, API_RAND, 1, HIP_UNSUPPORTED}},
-  {"curandDistributionShift_t",                     {"hiprandDistributionShift_t",                     "", CONV_TYPE, API_RAND, 1, HIP_UNSUPPORTED}},
+  {"curandDirectionVectorSet",                      {"hiprandDirectionVectorSet_t",                    "rocrand_direction_vector_set",                                   CONV_TYPE, API_RAND, 1}},
+  {"curandDirectionVectorSet_t",                    {"hiprandDirectionVectorSet_t",                    "rocrand_direction_vector_set",                                   CONV_TYPE, API_RAND, 1}},
+  // RAND choice of direction vector set (enum curandDirectionVectorSet)
+  {"CURAND_DIRECTION_VECTORS_32_JOEKUO6",           {"HIPRAND_DIRECTION_VECTORS_32_JOEKUO6",           "ROCRAND_DIRECTION_VECTORS_32_JOEKUO6",                           CONV_NUMERIC_LITERAL, API_RAND, 1}},
+  {"CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", {"HIPRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6", "ROCRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6",                 CONV_NUMERIC_LITERAL, API_RAND, 1}},
+  {"CURAND_DIRECTION_VECTORS_64_JOEKUO6",           {"HIPRAND_DIRECTION_VECTORS_64_JOEKUO6",           "ROCRAND_DIRECTION_VECTORS_64_JOEKUO6",                           CONV_NUMERIC_LITERAL, API_RAND, 1}},
+  {"CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", {"HIPRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6", "ROCRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6",                 CONV_NUMERIC_LITERAL, API_RAND, 1}},
+
+  {"curandGenerator_st",                            {"hiprandGenerator_st",                            "rocrand_generator_base_type",                                    CONV_TYPE, API_RAND, 1}},
+  {"curandGenerator_t",                             {"hiprandGenerator_t",                             "rocrand_generator",                                              CONV_TYPE, API_RAND, 1}},
+
+  {"curandDistribution_st",                         {"hiprandDistribution_st",                         "", CONV_TYPE, API_RAND, 1, UNSUPPORTED}},
+  {"curandDistribution_t",                          {"hiprandDistribution_t",                          "", CONV_TYPE, API_RAND, 1, UNSUPPORTED}},
+
+  {"curandHistogramM2V_st",                         {"hiprandHistogramM2V_st",                         "", CONV_TYPE, API_RAND, 1, UNSUPPORTED}},
+  {"curandHistogramM2V_t",                          {"hiprandHistogramM2V_t",                          "", CONV_TYPE, API_RAND, 1, UNSUPPORTED}},
+
+  {"curandDistributionShift_st",                    {"hiprandDistributionShift_st",                    "", CONV_TYPE, API_RAND, 1, UNSUPPORTED}},
+  {"curandDistributionShift_t",                     {"hiprandDistributionShift_t",                     "", CONV_TYPE, API_RAND, 1, UNSUPPORTED}},
+
   {"curandDistributionM2Shift_st",                  {"hiprandDistributionM2Shift_st",                  "", CONV_TYPE, API_RAND, 1, HIP_UNSUPPORTED}},
   {"curandDistributionM2Shift_t",                   {"hiprandDistributionM2Shift_t",                   "", CONV_TYPE, API_RAND, 1, HIP_UNSUPPORTED}},
   {"curandHistogramM2_st",                          {"hiprandHistogramM2_st",                          "", CONV_TYPE, API_RAND, 1, HIP_UNSUPPORTED}},
@@ -111,12 +121,6 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RAND_TYPE_NAME_MAP {
   {"curandState",                                   {"hiprandState",                                   "", CONV_TYPE, API_RAND, 1}},
   {"curandState_t",                                 {"hiprandState_t",                                 "", CONV_TYPE, API_RAND, 1}},
 
-  // RAND choice of direction vector set (enum curandDirectionVectorSet)
-  {"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}},
   {"CURAND_ITR",                                    {"HIPRAND_ITR",                                    "", CONV_NUMERIC_LITERAL, API_RAND, 1, HIP_UNSUPPORTED}},
@@ -236,4 +240,11 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_RAND_TYPE_NAME_VER_MAP {
   {"ROCRAND_ORDERING_PSEUDO_LEGACY",                {HIP_5050, HIP_0,    HIP_0   }},
   {"ROCRAND_ORDERING_PSEUDO_DYNAMIC",               {HIP_5050, HIP_0,    HIP_0   }},
   {"ROCRAND_ORDERING_QUASI_DEFAULT",                {HIP_5050, HIP_0,    HIP_0   }},
+  {"rocrand_direction_vector_set",                  {HIP_6000, HIP_0,    HIP_0   }},
+  {"ROCRAND_DIRECTION_VECTORS_32_JOEKUO6",          {HIP_6000, HIP_0,    HIP_0   }},
+  {"ROCRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6",{HIP_6000, HIP_0,    HIP_0   }},
+  {"ROCRAND_DIRECTION_VECTORS_64_JOEKUO6",          {HIP_6000, HIP_0,    HIP_0   }},
+  {"ROCRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6",{HIP_6000, HIP_0,    HIP_0   }},
+  {"rocrand_generator_base_type",                   {HIP_1050, HIP_0,    HIP_0   }},
+  {"rocrand_generator",                             {HIP_1050, HIP_0,    HIP_0   }},
 };
diff --git a/tests/unit_tests/synthetic/libraries/curand2hiprand.cu b/tests/unit_tests/synthetic/libraries/curand2hiprand.cu
@@ -120,8 +120,16 @@ int main() {
 
   // CHECK: hiprandDirectionVectorSet_t directionVectorSet;
   // CHECK-NEXT: hiprandDirectionVectorSet_t directionVectorSet_t;
+  // CHECK-NEXT: hiprandDirectionVectorSet_t DIRECTION_VECTORS_32_JOEKUO6 = HIPRAND_DIRECTION_VECTORS_32_JOEKUO6;
+  // CHECK-NEXT: hiprandDirectionVectorSet_t SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6 = HIPRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6;
+  // CHECK-NEXT: hiprandDirectionVectorSet_t DIRECTION_VECTORS_64_JOEKUO6 = HIPRAND_DIRECTION_VECTORS_64_JOEKUO6;
+  // CHECK-NEXT: hiprandDirectionVectorSet_t SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6 = HIPRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6;
   curandDirectionVectorSet directionVectorSet;
   curandDirectionVectorSet_t directionVectorSet_t;
+  curandDirectionVectorSet_t DIRECTION_VECTORS_32_JOEKUO6 = CURAND_DIRECTION_VECTORS_32_JOEKUO6;
+  curandDirectionVectorSet_t SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6 = CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6;
+  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 curandSetGeneratorOrdering(curandGenerator_t generator, curandOrdering_t order);
   // HIP: hiprandStatus_t HIPRANDAPI hiprandSetGeneratorOrdering(hiprandGenerator_t generator, hiprandOrdering_t order);

diff --git a/tests/unit_tests/synthetic/libraries/curand2rocrand.cu b/tests/unit_tests/synthetic/libraries/curand2rocrand.cu
@@ -74,6 +74,24 @@ int main() {
   curandOrdering_t RAND_ORDERING_PSEUDO_SEEDED = CURAND_ORDERING_PSEUDO_SEEDED;
   curandOrdering_t RAND_ORDERING_QUASI_DEFAULT = CURAND_ORDERING_QUASI_DEFAULT;
 
+  // CHECK: rocrand_direction_vector_set directionVectorSet;
+  // CHECK-NEXT: rocrand_direction_vector_set directionVectorSet_t;
+  // CHECK-NEXT: rocrand_direction_vector_set DIRECTION_VECTORS_32_JOEKUO6 = ROCRAND_DIRECTION_VECTORS_32_JOEKUO6;
+  // CHECK-NEXT: rocrand_direction_vector_set SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6 = ROCRAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6;
+  // CHECK-NEXT: rocrand_direction_vector_set DIRECTION_VECTORS_64_JOEKUO6 = ROCRAND_DIRECTION_VECTORS_64_JOEKUO6;
+  // CHECK-NEXT: rocrand_direction_vector_set SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6 = ROCRAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6;
+  curandDirectionVectorSet directionVectorSet;
+  curandDirectionVectorSet_t directionVectorSet_t;
+  curandDirectionVectorSet_t DIRECTION_VECTORS_32_JOEKUO6 = CURAND_DIRECTION_VECTORS_32_JOEKUO6;
+  curandDirectionVectorSet_t SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6 = CURAND_SCRAMBLED_DIRECTION_VECTORS_32_JOEKUO6;
+  curandDirectionVectorSet_t DIRECTION_VECTORS_64_JOEKUO6 = CURAND_DIRECTION_VECTORS_64_JOEKUO6;
+  curandDirectionVectorSet_t SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6 = CURAND_SCRAMBLED_DIRECTION_VECTORS_64_JOEKUO6;
+
+  // CHECK: rocrand_generator_base_type *randGenerator_st = nullptr;
+  // CHECK-NEXT: rocrand_generator randGenerator;
+  curandGenerator_st *randGenerator_st = nullptr;
+  curandGenerator_t randGenerator;
+
 #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;