diff --git a/clang/lib/DPCT/APINames_cuRAND.inc b/clang/lib/DPCT/APINames_cuRAND.inc index a1c6001e19ed..e5b3bc4eba3e 100644 --- a/clang/lib/DPCT/APINames_cuRAND.inc +++ b/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") diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index 2fb50f62b31b..2a63a1c52236 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/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(DE->getDecl()); emplaceTransformation(new ReplaceStmt(DE, toString(EC->getInitVal(), 10))); } + if (const DeclRefExpr *DE = + getNodeAsType(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(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, diff --git a/clang/lib/DPCT/ExprAnalysis.cpp b/clang/lib/DPCT/ExprAnalysis.cpp index ede69a66c887..8815eca84169 100644 --- a/clang/lib/DPCT/ExprAnalysis.cpp +++ b/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 diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index b40f771f4d71..1f6eb329f82e 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -33,6 +33,7 @@ std::unordered_map> std::unordered_map> MapNames::ClassFieldMap; MapNames::MapTy MapNames::RandomEngineTypeMap; +MapNames::MapTy MapNames::RandomOrderingTypeMap; MapNames::MapTy MapNames::DeviceRandomGeneratorTypeMap; std::unordered_map> MapNames::CuDNNTypeNamesMap; @@ -343,6 +344,7 @@ void MapNames::setExplicitNamespaceMap() { HelperFeatureEnum::device_ext)}, {"curandStatus_t", std::make_shared("int")}, {"curandStatus", std::make_shared("int")}, + {"curandOrdering_t", std::make_shared("uint32_t")}, {"cusparseStatus_t", std::make_shared("int")}, {"cusparseMatDescr_t", std::make_shared("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() + diff --git a/clang/lib/DPCT/MapNames.h b/clang/lib/DPCT/MapNames.h index 54c3065a0583..3e566aa18e03 100644 --- a/clang/lib/DPCT/MapNames.h +++ b/clang/lib/DPCT/MapNames.h @@ -356,6 +356,7 @@ class MapNames { static MapTy ITFName; static MapTy RandomEngineTypeMap; + static MapTy RandomOrderingTypeMap; static const std::map RandomGenerateFuncMap; static MapTy DeviceRandomGeneratorTypeMap; diff --git a/clang/runtime/dpct-rt/include/dpct/rng_utils.hpp b/clang/runtime/dpct-rt/include/dpct/rng_utils.hpp index 306b35053068..3b1dcf41420c 100644 --- a/clang/runtime/dpct-rt/include/dpct/rng_utils.hpp +++ b/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 _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) { + 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(*queue, dimensions) : engine_t(*queue, seed); diff --git a/clang/test/dpct/curand.cu b/clang/test/dpct/curand.cu index 536b9a516f5f..54692c5c156a 100644 --- a/clang/test/dpct/curand.cu +++ b/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); diff --git a/clang/test/dpct/curandEnum.cu b/clang/test/dpct/curandEnum.cu index 68c27d4cbec8..f041c795b78a 100644 --- a/clang/test/dpct/curandEnum.cu +++ b/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 + ); +}