[SYCLomatic] Support migration of curandSetGeneratorOrdering.

Signed-off-by: Tang, Jiajun [email protected]

clang/lib/DPCT/APINames_cuRAND.inc

@@ -53,7 +53,7 @@ ENTRY(curandGetScrambleConstants32, curandGetScrambleConstants32, false, NO_FLAG
 ENTRY(curandGetScrambleConstants64, curandGetScrambleConstants64, false, NO_FLAG, P4, "comment")
 ENTRY(curandGetVersion, curandGetVersion, false, NO_FLAG, P4, "comment")
 ENTRY(curandSetGeneratorOffset, curandSetGeneratorOffset, true, NO_FLAG, P4, "Successful")
-ENTRY(curandSetGeneratorOrdering, curandSetGeneratorOrdering, false, NO_FLAG, P4, "comment")
+ENTRY(curandSetGeneratorOrdering, curandSetGeneratorOrdering, true, NO_FLAG, P4, "Successful")
 ENTRY(curandSetPseudoRandomGeneratorSeed, curandSetPseudoRandomGeneratorSeed, true, NO_FLAG, P4, "Successful")
 ENTRY(curandSetQuasiRandomGeneratorDimensions, curandSetQuasiRandomGeneratorDimensions, true, NO_FLAG, P4, "Successful")
 ENTRY(curandSetStream, curandSetStream, true, NO_FLAG, P4, "Successful")

clang/lib/DPCT/ASTTraversal.cpp

@@ -1722,17 +1722,18 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
               "cublasOperation_t", "cusolverStatus_t", "cusolverEigType_t",
               "cusolverEigMode_t", "curandStatus_t", "cudaStream_t",
               "cusparseStatus_t", "cusparseDiagType_t", "cusparseFillMode_t",
-              "cusparseIndexBase_t", "cusparseMatrixType_t", "cusparseAlgMode_t",
-              "cusparseOperation_t", "cusparseMatDescr_t", "cusparseHandle_t",
-              "CUcontext", "cublasPointerMode_t", "cusparsePointerMode_t",
-              "cublasGemmAlgo_t", "cusparseSolveAnalysisInfo_t", "cudaDataType",
-              "cublasDataType_t", "curandState_t", "curandState",
-              "curandStateXORWOW_t", "curandStateXORWOW",
-              "curandStatePhilox4_32_10_t", "curandStatePhilox4_32_10",
-              "curandStateMRG32k3a_t", "curandStateMRG32k3a", "thrust::minus",
-              "thrust::negate", "thrust::logical_or", "thrust::equal_to",
-              "thrust::less", "cudaSharedMemConfig", "curandGenerator_t",
-              "curandRngType_t", "cufftHandle", "cufftReal", "cufftDoubleReal",
+              "cusparseIndexBase_t", "cusparseMatrixType_t",
+              "cusparseAlgMode_t", "cusparseOperation_t", "cusparseMatDescr_t",
+              "cusparseHandle_t", "CUcontext", "cublasPointerMode_t",
+              "cusparsePointerMode_t", "cublasGemmAlgo_t",
+              "cusparseSolveAnalysisInfo_t", "cudaDataType", "cublasDataType_t",
+              "curandState_t", "curandState", "curandStateXORWOW_t",
+              "curandStateXORWOW", "curandStatePhilox4_32_10_t",
+              "curandStatePhilox4_32_10", "curandStateMRG32k3a_t",
+              "curandStateMRG32k3a", "thrust::minus", "thrust::negate",
+              "thrust::logical_or", "thrust::equal_to", "thrust::less",
+              "cudaSharedMemConfig", "curandGenerator_t", "curandRngType_t",
+              "curandOrdering_t", "cufftHandle", "cufftReal", "cufftDoubleReal",
               "cufftComplex", "cufftDoubleComplex", "cufftResult_t",
               "cufftResult", "cufftType_t", "cufftType", "thrust::pair",
               "CUdeviceptr", "cudaDeviceAttr", "CUmodule", "CUjit_option",
@@ -1750,7 +1751,8 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
               "cusparseConstDnVecDescr_t", "cusparseSpMatDescr_t",
               "cusparseSpMMAlg_t", "cusparseSpMVAlg_t", "cusparseSpGEMMDescr_t",
               "cusparseSpSVDescr_t", "cusparseSpGEMMAlg_t", "CUuuid",
-              "cusparseSpSVAlg_t", "cudaFuncAttributes", "cudaLaunchAttributeValue"))))))
+              "cusparseSpSVAlg_t", "cudaFuncAttributes",
+              "cudaLaunchAttributeValue"))))))
           .bind("cudaTypeDef"),
       this);
   MF.addMatcher(varDecl(hasType(classTemplateSpecializationDecl(
@@ -3656,6 +3658,10 @@ void RandomEnumsRule::registerMatcher(MatchFinder &MF) {
       declRefExpr(to(enumConstantDecl(matchesName("CURAND_STATUS.*"))))
           .bind("RANDOMStatusConstants"),
       this);
+  MF.addMatcher(
+      declRefExpr(to(enumConstantDecl(matchesName("CURAND_ORDERING.*"))))
+          .bind("RANDOMOrderingConstants"),
+      this);
   MF.addMatcher(declRefExpr(to(enumConstantDecl(matchesName("CURAND_RNG.*"))))
                     .bind("RandomTypeEnum"),
                 this);
@@ -3667,6 +3673,16 @@ void RandomEnumsRule::runRule(const MatchFinder::MatchResult &Result) {
     auto *EC = cast<EnumConstantDecl>(DE->getDecl());
     emplaceTransformation(new ReplaceStmt(DE, toString(EC->getInitVal(), 10)));
   }
+  if (const DeclRefExpr *DE =
+          getNodeAsType<DeclRefExpr>(Result, "RANDOMOrderingConstants")) {
+    std::string EnumStr = DE->getNameInfo().getName().getAsString();
+    auto Search = MapNames::RandomOrderingTypeMap.find(EnumStr);
+    if (Search == MapNames::RandomOrderingTypeMap.end()) {
+      report(DE->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false, EnumStr);
+      return;
+    }
+    emplaceTransformation(new ReplaceStmt(DE, Search->second));
+  }
   if (const DeclRefExpr *DE =
           getNodeAsType<DeclRefExpr>(Result, "RandomTypeEnum")) {
     std::string EnumStr = DE->getNameInfo().getName().getAsString();
@@ -3888,7 +3904,7 @@ void RandomFunctionCallRule::registerMatcher(MatchFinder &MF) {
         "curandGenerateNormal", "curandGenerateNormalDouble",
         "curandGeneratePoisson", "curandGenerateUniform",
         "curandGenerateUniformDouble", "curandSetStream",
-        "curandCreateGeneratorHost");
+        "curandCreateGeneratorHost", "curandSetGeneratorOrdering");
   };
   MF.addMatcher(
       callExpr(allOf(callee(functionDecl(functionName())), parentStmt()))
@@ -4013,6 +4029,12 @@ void RandomFunctionCallRule::runRule(const MatchFinder::MatchResult &Result) {
         buildString(ExprAnalysis::ref(CE->getArg(0)), "->set_queue(",
                     ExprAnalysis::ref(CE->getArg(1)), ")")));
   }
+  if (FuncName == "curandSetGeneratorOrdering") {
+    return emplaceTransformation(new ReplaceStmt(
+        CE, false,
+        buildString(ExprAnalysis::ref(CE->getArg(0)), "->set_custom(",
+                    ExprAnalysis::ref(CE->getArg(1)), ")")));
+  }
 }
 
 REGISTER_RULE(RandomFunctionCallRule, PassKind::PK_Migration,

clang/lib/DPCT/ExprAnalysis.cpp

@@ -564,6 +564,7 @@ void ExprAnalysis::analyzeExpr(const DeclRefExpr *DRE) {
       REPLACE_ENUM(MapNames::FunctionAttrMap);
       REPLACE_ENUM(CuDNNTypeRule::CuDNNEnumNamesMap);
       REPLACE_ENUM(MapNames::RandomEngineTypeMap);
+      REPLACE_ENUM(MapNames::RandomOrderingTypeMap);
       REPLACE_ENUM(MapNames::SOLVEREnumsMap);
       REPLACE_ENUM(MapNames::SPBLASEnumsMap);
 #undef REPLACE_ENUM

clang/lib/DPCT/MapNames.cpp

@@ -33,6 +33,7 @@ std::unordered_map<std::string, std::shared_ptr<TypeNameRule>>
 std::unordered_map<std::string, std::shared_ptr<ClassFieldRule>>
     MapNames::ClassFieldMap;
 MapNames::MapTy MapNames::RandomEngineTypeMap;
+MapNames::MapTy MapNames::RandomOrderingTypeMap;
 MapNames::MapTy MapNames::DeviceRandomGeneratorTypeMap;
 std::unordered_map<std::string, std::shared_ptr<TypeNameRule>>
     MapNames::CuDNNTypeNamesMap;
@@ -343,6 +344,7 @@ void MapNames::setExplicitNamespaceMap() {
                             HelperFeatureEnum::device_ext)},
       {"curandStatus_t", std::make_shared<TypeNameRule>("int")},
       {"curandStatus", std::make_shared<TypeNameRule>("int")},
+      {"curandOrdering_t", std::make_shared<TypeNameRule>("uint32_t")},
       {"cusparseStatus_t", std::make_shared<TypeNameRule>("int")},
       {"cusparseMatDescr_t",
        std::make_shared<TypeNameRule>("std::shared_ptr<" + getDpctNamespace() +
@@ -514,6 +516,16 @@ void MapNames::setExplicitNamespaceMap() {
        getDpctNamespace() + "rng::random_engine_type::sobol"},
   };
 
+  // Random Ordering Type mapping
+  RandomOrderingTypeMap = {
+      {"CURAND_ORDERING_PSEUDO_DEFAULT", "81920"},
+      {"CURAND_ORDERING_PSEUDO_BEST", "81920"},
+      // CURAND_ORDERING_PSEUDO_SEEDED not support now.
+      {"CURAND_ORDERING_PSEUDO_LEGACY", "4096"},
+      {"CURAND_ORDERING_PSEUDO_DYNAMIC", "0"},
+      // CURAND_ORDERING_QUASI_DEFAULT not support now.
+  };
+
   // Device Random Generator Type mapping
   DeviceRandomGeneratorTypeMap = {
       {"curandStateXORWOW_t", getDpctNamespace() +

clang/lib/DPCT/MapNames.h

@@ -356,6 +356,7 @@ class MapNames {
 
   static MapTy ITFName;
   static MapTy RandomEngineTypeMap;
+  static MapTy RandomOrderingTypeMap;
   static const std::map<std::string, std::string> RandomGenerateFuncMap;
 
   static MapTy DeviceRandomGeneratorTypeMap;

clang/runtime/dpct-rt/include/dpct/rng_utils.hpp

@@ -221,6 +221,10 @@ class rng_generator_base {
   /// \param queue The engine queue.
   virtual void set_queue(sycl::queue *queue) = 0;
 
+  /// Set the custom of host rng_generator.
+  /// \param custom The engine custom.
+  virtual void set_custom(const std::uint32_t custom) = 0;
+
   /// Generate unsigned int random number(s) with 'uniform_bits' distribution.
   /// \param output The pointer of the first random number.
   /// \param n The number of random numbers.
@@ -301,6 +305,7 @@ class rng_generator_base {
   sycl::queue *_queue = nullptr;
   std::uint64_t _seed{0};
   std::uint32_t _dimensions{1};
+  std::uint32_t _custom{81920};
   std::vector<std::uint32_t> _direction_numbers;
 };
 
@@ -311,8 +316,8 @@ class rng_generator : public rng_generator_base {
   /// Constructor of rng_generator.
   /// \param q The queue where the generator should be executed.
   rng_generator(sycl::queue &q = dpct::get_default_queue())
-      : rng_generator_base(&q), _engine(create_engine(&q, _seed, _dimensions)) {
-  }
+      : rng_generator_base(&q),
+        _engine(create_engine(&q, _seed, _dimensions, _custom)) {}
 
   /// Set the seed of host rng_generator.
   /// \param seed The engine seed.
@@ -321,7 +326,7 @@ class rng_generator : public rng_generator_base {
       return;
     }
     _seed = seed;
-    _engine = create_engine(_queue, _seed, _dimensions);
+    _engine = create_engine(_queue, _seed, _dimensions, _custom);
   }
 
   /// Set the dimensions of host rng_generator.
@@ -331,7 +336,7 @@ class rng_generator : public rng_generator_base {
       return;
     }
     _dimensions = dimensions;
-    _engine = create_engine(_queue, _seed, _dimensions);
+    _engine = create_engine(_queue, _seed, _dimensions, _custom);
   }
 
   /// Set the queue of host rng_generator.
@@ -341,7 +346,22 @@ class rng_generator : public rng_generator_base {
       return;
     }
     _queue = queue;
-    _engine = create_engine(_queue, _seed, _dimensions);
+    _engine = create_engine(_queue, _seed, _dimensions, _custom);
+  }
+
+  /// Set the custom of host rng_generator.
+  /// \param custom The engine custom.
+  void set_custom(const std::uint32_t custom) {
+#ifndef __INTEL_MKL__
+    throw std::runtime_error("The oneAPI Math Kernel Library (oneMKL) "
+                             "Interfaces Project does not support this API.");
+#else
+    if (custom == _custom) {
+      return;
+    }
+    _custom = custom;
+    _engine = create_engine(_queue, _seed, _dimensions, _custom);
+#endif
   }
 
   /// Set the direction numbers of Sobol host rng_generator.
@@ -473,8 +493,15 @@ class rng_generator : public rng_generator_base {
 private:
   static inline engine_t create_engine(sycl::queue *queue,
                                        const std::uint64_t seed,
-                                       const std::uint32_t dimensions) {
+                                       const std::uint32_t dimensions,
+                                       const std::uint32_t custom) {
 #ifdef __INTEL_MKL__
+    if constexpr (std::is_same_v<engine_t, oneapi::mkl::rng::mrg32k3a>) {
+      if (custom)
+        return engine_t(*queue, seed,
+                        oneapi::mkl::rng::mrg32k3a_mode::custom{custom});
+      return engine_t(*queue, seed, oneapi::mkl::rng::mrg32k3a_mode::optimal_v);
+    }
     return std::is_same_v<engine_t, oneapi::mkl::rng::sobol>
                ? engine_t(*queue, dimensions)
                : engine_t(*queue, seed);

clang/test/dpct/curand.cu

@@ -79,6 +79,11 @@ int main(){
   curandSetGeneratorOffset(rng, 100);
   s1 = curandSetGeneratorOffset(rng2, 200);
 
+  //CHECK:rng->set_custom(81920);
+  //CHECK-NEXT:s1 = DPCT_CHECK_ERROR(rng2->set_custom(81920));
+  curandSetGeneratorOrdering(rng, CURAND_ORDERING_PSEUDO_BEST);
+  s1 = curandSetGeneratorOrdering(rng2, CURAND_ORDERING_PSEUDO_DEFAULT);
+
   //CHECK:rng.reset();
   //CHECK-NEXT:s1 = DPCT_CHECK_ERROR(rng.reset());
   curandDestroyGenerator(rng);

clang/test/dpct/curandEnum.cu

@@ -1,3 +1,5 @@
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2
+// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2
 //RUN: dpct --out-root %T/curandEnum --format-range=none --cuda-include-path="%cuda-path/include" %s -- -x cuda --cuda-host-only
 //RUN: FileCheck --input-file %T/curandEnum/curandEnum.dp.cpp --match-full-lines %s
 // RUN: %if build_lit %{icpx -c -fsycl %T/curandEnum/curandEnum.dp.cpp -o %T/curandEnum/curandEnum.dp.o %}
@@ -35,6 +37,23 @@ curandStatus_t foo(
   curandStatus_t a12,
   curandStatus_t a13) {}
 
+// CHECK:uint32_t goo(
+// CHECK-NEXT:uint32_t b1,
+// CHECK-NEXT:uint32_t b2,
+// CHECK-NEXT:// curandOrdering_t b3,
+// CHECK-NEXT:uint32_t b4,
+// CHECK-NEXT:uint32_t b5
+// CHECK-NEXT:// , curandOrdering_t b6
+// CHECK-NEXT:) { return b1; }
+curandOrdering_t goo(
+    curandOrdering_t b1,
+    curandOrdering_t b2,
+    // curandOrdering_t b3,
+    curandOrdering_t b4,
+    curandOrdering_t b5
+    // , curandOrdering_t b6
+) { return b1; }
+
 int main() {
   // CHECK:int a1 = 0;
   // CHECK-NEXT:int a2 = 100;
@@ -92,5 +111,34 @@ int main() {
     CURAND_STATUS_INITIALIZATION_FAILED,
     CURAND_STATUS_ARCH_MISMATCH,
     CURAND_STATUS_INTERNAL_ERROR);
-}
 
+  // CHECK:uint32_t b1 = 81920;
+  // CHECK-NEXT:uint32_t b2 = 81920;
+  // CHECK-NEXT:// curandOrdering_t b3 = CURAND_ORDERING_PSEUDO_SEEDED;
+  // CHECK-NEXT:uint32_t b4 = 4096;
+  // CHECK-NEXT:uint32_t b5 = 0;
+  // CHECK-NEXT:// curandOrdering_t b6 = CURAND_ORDERING_QUASI_DEFAULT;
+  curandOrdering_t b1 = CURAND_ORDERING_PSEUDO_BEST;
+  curandOrdering_t b2 = CURAND_ORDERING_PSEUDO_DEFAULT;
+  // curandOrdering_t b3 = CURAND_ORDERING_PSEUDO_SEEDED;
+  curandOrdering_t b4 = CURAND_ORDERING_PSEUDO_LEGACY;
+  curandOrdering_t b5 = CURAND_ORDERING_PSEUDO_DYNAMIC;
+  // curandOrdering_t b6 = CURAND_ORDERING_QUASI_DEFAULT;
+
+  // CHECK:goo(
+  // CHECK-NEXT:  81920,
+  // CHECK-NEXT:  81920,
+  // CHECK-NEXT:  // CURAND_ORDERING_PSEUDO_SEEDED,
+  // CHECK-NEXT:  4096,
+  // CHECK-NEXT:  0
+  // CHECK-NEXT:  // , CURAND_ORDERING_QUASI_DEFAULT
+  // CHECK-NEXT:);
+  goo(
+      CURAND_ORDERING_PSEUDO_BEST,
+      CURAND_ORDERING_PSEUDO_DEFAULT,
+      // CURAND_ORDERING_PSEUDO_SEEDED,
+      CURAND_ORDERING_PSEUDO_LEGACY,
+      CURAND_ORDERING_PSEUDO_DYNAMIC
+      // , CURAND_ORDERING_QUASI_DEFAULT
+  );
+}