diff --git a/docs/markdown/CURAND_API_supported_by_HIP.md b/docs/markdown/CURAND_API_supported_by_HIP.md index 900cfadc03..aae43af272 100644 --- a/docs/markdown/CURAND_API_supported_by_HIP.md +++ b/docs/markdown/CURAND_API_supported_by_HIP.md @@ -154,6 +154,8 @@ | `curand_normal2_double` | `hiprand_normal2_double` | | `curand_normal4` | `hiprand_normal4` | | `curand_normal4_double` | `hiprand_normal4_double` | +| `curand_uniform` | `hiprand_uniform` | +| `curand_uniform_double` | `hiprand_uniform_double` | | `curand_uniform2_double` | `hiprand_uniform2_double` | | `curand_uniform4` | `hiprand_uniform4` | | `curand_uniform4_double` | `hiprand_uniform4_double` | diff --git a/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp index 6bcc0b38c4..7001deb2e7 100644 --- a/hipify-clang/src/CUDA2HipMap.cpp +++ b/hipify-clang/src/CUDA2HipMap.cpp @@ -367,24 +367,39 @@ const std::map CUDA_TYPE_NAME_MAP{ /// Maps cuda header names to hip header names. const std::map CUDA_INCLUDE_MAP{ // CUDA includes - {"cuda.h", {"hip/hip_runtime.h", CONV_INCLUDE_CUDA_MAIN_H, API_DRIVER}}, - {"cuda_runtime.h", {"hip/hip_runtime.h", CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME}}, - {"cuda_runtime_api.h", {"hip/hip_runtime_api.h", CONV_INCLUDE, API_RUNTIME}}, - {"channel_descriptor.h", {"hip/channel_descriptor.h", CONV_INCLUDE, API_RUNTIME}}, - {"device_functions.h", {"hip/device_functions.h", CONV_INCLUDE, API_RUNTIME}}, - {"driver_types.h", {"hip/driver_types.h", CONV_INCLUDE, API_RUNTIME}}, - {"cuComplex.h", {"hip/hip_complex.h", CONV_INCLUDE, API_RUNTIME}}, - {"cuda_fp16.h", {"hip/hip_fp16.h", CONV_INCLUDE, API_RUNTIME}}, - {"cuda_texture_types.h", {"hip/hip_texture_types.h", CONV_INCLUDE, API_RUNTIME}}, - {"vector_types.h", {"hip/hip_vector_types.h", CONV_INCLUDE, API_RUNTIME}}, + {"cuda.h", {"hip/hip_runtime.h", CONV_INCLUDE_CUDA_MAIN_H, API_DRIVER}}, + {"cuda_runtime.h", {"hip/hip_runtime.h", CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME}}, + {"cuda_runtime_api.h", {"hip/hip_runtime_api.h", CONV_INCLUDE, API_RUNTIME}}, + {"channel_descriptor.h", {"hip/channel_descriptor.h", CONV_INCLUDE, API_RUNTIME}}, + {"device_functions.h", {"hip/device_functions.h", CONV_INCLUDE, API_RUNTIME}}, + {"driver_types.h", {"hip/driver_types.h", CONV_INCLUDE, API_RUNTIME}}, + {"cuComplex.h", {"hip/hip_complex.h", CONV_INCLUDE, API_RUNTIME}}, + {"cuda_fp16.h", {"hip/hip_fp16.h", CONV_INCLUDE, API_RUNTIME}}, + {"cuda_texture_types.h", {"hip/hip_texture_types.h", CONV_INCLUDE, API_RUNTIME}}, + {"vector_types.h", {"hip/hip_vector_types.h", CONV_INCLUDE, API_RUNTIME}}, // CUBLAS includes - {"cublas.h", {"hipblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}}, - {"cublas_v2.h", {"hipblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}}, + {"cublas.h", {"hipblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}}, + {"cublas_v2.h", {"hipblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}}, // CURAND includes - {"curand.h", {"hiprand.h", CONV_INCLUDE, API_RAND}}, - {"curand_kernel.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand.h", {"hiprand.h", CONV_INCLUDE_CUDA_MAIN_H, API_RAND}}, + {"curand_kernel.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_discrete.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_discrete2.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_globals.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_lognormal.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_mrg32k3a.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_mtgp32.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_mtgp32_host.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_mtgp32_kernel.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_mtgp32dc_p_11213.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_normal.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_normal_static.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_philox4x32_x.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_poisson.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_precalc.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand_uniform.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, // HIP includes // TODO: uncomment this when hip/cudacommon.h will be renamed to hip/hipcommon.h @@ -2852,6 +2867,8 @@ const std::map CUDA_IDENTIFIER_MAP{ {"curand_normal2_double", {"hiprand_normal2_double", CONV_DEVICE_FUNC, API_RAND}}, {"curand_normal4", {"hiprand_normal4", CONV_DEVICE_FUNC, API_RAND}}, {"curand_normal4_double", {"hiprand_normal4_double", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_uniform", {"hiprand_uniform", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_uniform_double", {"hiprand_uniform_double", CONV_DEVICE_FUNC, API_RAND}}, {"curand_uniform2_double", {"hiprand_uniform2_double", CONV_DEVICE_FUNC, API_RAND}}, {"curand_uniform4", {"hiprand_uniform4", CONV_DEVICE_FUNC, API_RAND}}, {"curand_uniform4_double", {"hiprand_uniform4_double", CONV_DEVICE_FUNC, API_RAND}}, diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index 3459ea21f9..76efe9bf85 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/hipify-clang/src/HipifyAction.cpp @@ -137,6 +137,48 @@ std::string stringifyZeroDefaultedArg(clang::SourceManager& SM, const clang::Exp } // anonymous namespace +bool HipifyAction::Exclude(const hipCounter & hipToken) { + switch (hipToken.type) { + case CONV_INCLUDE_CUDA_MAIN_H: + switch (hipToken.apiType) { + case API_DRIVER: + case API_RUNTIME: + if (insertedRuntimeHeader) { return true; } + insertedRuntimeHeader = true; + return false; + case API_BLAS: + if (insertedBLASHeader) { return true; } + insertedBLASHeader = true; + return false; + case API_RAND: + if (hipToken.hipName == "hiprand_kernel.h") { + if (insertedRAND_kernelHeader) { return true; } + insertedRAND_kernelHeader = true; + return false; + } else if (hipToken.hipName == "hiprand.h") { + if (insertedRANDHeader) { return true; } + insertedRANDHeader = true; + return false; + } + default: + return false; + } + return false; + case CONV_INCLUDE: + switch (hipToken.apiType) { + case API_RAND: + if (insertedRAND_kernelHeader) { return true; } + insertedRAND_kernelHeader = true; + return false; + default: + return false; + } + return false; + default: + return false; + } + return false; +} void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, const clang::Token&, @@ -159,29 +201,7 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, return; } - // Special-casing to avoid duplication of the hip_runtime include. - bool secondMainInclude = false; - if (found->second.countType == CONV_INCLUDE_CUDA_MAIN_H) { - switch (found->second.countApiType) { - case API_DRIVER: - case API_RUNTIME: - if (insertedRuntimeHeader) { - secondMainInclude = true; - break; - } - insertedRuntimeHeader = true; - break; - case API_BLAS: - if (insertedBLASHeader) { - secondMainInclude = true; - break; - } - insertedBLASHeader = true; - break; - default: - break; - } - } + bool exclude = Exclude(found->second); Statistics::current().incrementCounter(found->second, file_name.str()); @@ -195,7 +215,7 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, clang::StringRef newInclude; // Keep the same include type that the user gave. - if (!secondMainInclude) { + if (!exclude) { clang::SmallString<128> includeBuffer; if (is_angled) { newInclude = llvm::Twine("<" + found->second.hipName + ">").toStringRef(includeBuffer); diff --git a/hipify-clang/src/HipifyAction.h b/hipify-clang/src/HipifyAction.h index 42622c1e01..8c2195b28d 100644 --- a/hipify-clang/src/HipifyAction.h +++ b/hipify-clang/src/HipifyAction.h @@ -6,6 +6,7 @@ #include "clang/Tooling/Core/Replacement.h" #include "clang/ASTMatchers/ASTMatchFinder.h" #include "ReplacementsFrontendActionFactory.h" +#include "Statistics.h" namespace ct = clang::tooling; @@ -24,6 +25,8 @@ private: // This approach means we do the best it's possible to do w.r.t preserving the user's include order. bool insertedRuntimeHeader = false; bool insertedBLASHeader = false; + bool insertedRANDHeader = false; + bool insertedRAND_kernelHeader = false; bool firstHeader = false; bool pragmaOnce = false; clang::SourceLocation firstHeaderLoc; @@ -90,4 +93,6 @@ protected: void run(const clang::ast_matchers::MatchFinder::MatchResult& Result) override; std::unique_ptr CreateASTConsumer(clang::CompilerInstance &CI, llvm::StringRef InFile) override; + + bool Exclude(const hipCounter & hipToken); }; diff --git a/hipify-clang/src/Statistics.cpp b/hipify-clang/src/Statistics.cpp index e2b3632066..4c5664a6b0 100644 --- a/hipify-clang/src/Statistics.cpp +++ b/hipify-clang/src/Statistics.cpp @@ -53,8 +53,8 @@ void printStat(std::ostream *csv, llvm::raw_ostream* printOut, const std::string void StatCounter::incrementCounter(const hipCounter& counter, std::string name) { counters[name]++; - apiCounters[(int) counter.countApiType]++; - convTypeCounters[(int) counter.countType]++; + apiCounters[(int) counter.apiType]++; + convTypeCounters[(int) counter.type]++; } void StatCounter::add(const StatCounter& other) { diff --git a/hipify-clang/src/Statistics.h b/hipify-clang/src/Statistics.h index f160ca7383..81be7b09a8 100644 --- a/hipify-clang/src/Statistics.h +++ b/hipify-clang/src/Statistics.h @@ -67,8 +67,8 @@ extern const char *apiNames[NUM_API_TYPES]; struct hipCounter { llvm::StringRef hipName; - ConvTypes countType; - ApiTypes countApiType; + ConvTypes type; + ApiTypes apiType; bool unsupported; }; diff --git a/tests/hipify-clang/cuRAND/benchmark_curand_generate.cpp b/tests/hipify-clang/cuRAND/benchmark_curand_generate.cpp new file mode 100644 index 0000000000..8d284c82c6 --- /dev/null +++ b/tests/hipify-clang/cuRAND/benchmark_curand_generate.cpp @@ -0,0 +1,393 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// Copyright (c) 2017 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#include +#include +#include +#include +#include +#include +#include +#include + +#include "cmdparser.hpp" +// CHECK: #include +#include +// CHECK: #include +#include + +// CHECK: if((x)!=hipSuccess) { +#define CUDA_CALL(x) do { if((x)!=cudaSuccess) { \ + printf("Error at %s:%d\n",__FILE__,__LINE__);\ + exit(EXIT_FAILURE);}} while(0) +// CHECK: if((x)!=HIPRAND_STATUS_SUCCESS) { +#define CURAND_CALL(x) do { if((x)!=CURAND_STATUS_SUCCESS) { \ + printf("Error at %s:%d\n",__FILE__,__LINE__);\ + exit(EXIT_FAILURE);}} while(0) + +#ifndef DEFAULT_RAND_N +const size_t DEFAULT_RAND_N = 1024 * 1024 * 128; +#endif + +// CHECK: typedef hiprandRngType_t rng_type_t; +typedef curandRngType rng_type_t; + +// CHECK: using generate_func_type = std::function; +template +using generate_func_type = std::function; + +template +void run_benchmark(const cli::Parser& parser, + const rng_type_t rng_type, + generate_func_type generate_func) +{ + const size_t size = parser.get("size"); + const size_t trials = parser.get("trials"); + + T * data; + // CHECK: CUDA_CALL(hipMalloc((void **)&data, size * sizeof(T))); + CUDA_CALL(cudaMalloc((void **)&data, size * sizeof(T))); + + // CHECK: hiprandGenerator_t generator; + // CHECK: CURAND_CALL(hiprandCreateGenerator(&generator, rng_type)); + curandGenerator_t generator; + CURAND_CALL(curandCreateGenerator(&generator, rng_type)); + + const size_t dimensions = parser.get("dimensions"); + // CHECK: hiprandStatus_t status = hiprandSetQuasiRandomGeneratorDimensions(generator, dimensions); + // CHECK: if (status != HIPRAND_STATUS_TYPE_ERROR) + curandStatus_t status = curandSetQuasiRandomGeneratorDimensions(generator, dimensions); + if (status != CURAND_STATUS_TYPE_ERROR) // If the RNG is not quasi-random + { + CURAND_CALL(status); + } + + // Warm-up + for (size_t i = 0; i < 5; i++) + { + CURAND_CALL(generate_func(generator, data, size)); + } + // CHECK: CUDA_CALL(hipDeviceSynchronize()); + CUDA_CALL(cudaDeviceSynchronize()); + + // Measurement + auto start = std::chrono::high_resolution_clock::now(); + for (size_t i = 0; i < trials; i++) + { + CURAND_CALL(generate_func(generator, data, size)); + } + // CHECK: CUDA_CALL(hipDeviceSynchronize()); + CUDA_CALL(cudaDeviceSynchronize()); + auto end = std::chrono::high_resolution_clock::now(); + std::chrono::duration elapsed = end - start; + + std::cout << std::fixed << std::setprecision(3) + << " " + << "Throughput = " + << std::setw(8) << (trials * size * sizeof(T)) / + (elapsed.count() / 1e3 * (1 << 30)) + << " GB/s, Samples = " + << std::setw(8) << (trials * size) / + (elapsed.count() / 1e3 * (1 << 30)) + << " GSample/s, AvgTime (1 trial) = " + << std::setw(8) << elapsed.count() / trials + << " ms, Time (all) = " + << std::setw(8) << elapsed.count() + << " ms, Size = " << size + << std::endl; + // CHECK: CURAND_CALL(hiprandDestroyGenerator(generator)); + // CHECK: CUDA_CALL(hipFree(data)); + CURAND_CALL(curandDestroyGenerator(generator)); + CUDA_CALL(cudaFree(data)); +} + +void run_benchmarks(const cli::Parser& parser, + const rng_type_t rng_type, + const std::string& distribution) +{ + if (distribution == "uniform-uint") + { + // CHECK: if (rng_type != HIPRAND_RNG_QUASI_SOBOL64 && + // CHECK: rng_type != HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64) + if (rng_type != CURAND_RNG_QUASI_SOBOL64 && + rng_type != CURAND_RNG_QUASI_SCRAMBLED_SOBOL64) + { + run_benchmark(parser, rng_type, + // CHECK: [](hiprandGenerator_t gen, unsigned int * data, size_t size) { + // CHECK: return hiprandGenerate(gen, data, size); + [](curandGenerator_t gen, unsigned int * data, size_t size) { + return curandGenerate(gen, data, size); + } + ); + } + } + if (distribution == "uniform-long-long") + { + // CHECK: if (rng_type == HIPRAND_RNG_QUASI_SOBOL64 || + // CHECK: rng_type == HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64) + if (rng_type == CURAND_RNG_QUASI_SOBOL64 || + rng_type == CURAND_RNG_QUASI_SCRAMBLED_SOBOL64) + { + run_benchmark(parser, rng_type, + // CHECK: [](hiprandGenerator_t gen, unsigned long long * data, size_t size) { + [](curandGenerator_t gen, unsigned long long * data, size_t size) { + // curandGenerateLongLong is yet unsupported by HIP + // CHECK-NOT: return hiprandGenerateLongLong(gen, data, size); + return curandGenerateLongLong(gen, data, size); + } + ); + } + } + if (distribution == "uniform-float") + { + run_benchmark(parser, rng_type, + // CHECK: [](hiprandGenerator_t gen, float * data, size_t size) { + // CHECK: return hiprandGenerateUniform(gen, data, size); + [](curandGenerator_t gen, float * data, size_t size) { + return curandGenerateUniform(gen, data, size); + } + ); + } + if (distribution == "uniform-double") + { + run_benchmark(parser, rng_type, + // CHECK: [](hiprandGenerator_t gen, double * data, size_t size) { + // CHECK: return hiprandGenerateUniformDouble(gen, data, size); + [](curandGenerator_t gen, double * data, size_t size) { + return curandGenerateUniformDouble(gen, data, size); + } + ); + } + if (distribution == "normal-float") + { + run_benchmark(parser, rng_type, + // CHECK: [](hiprandGenerator_t gen, float * data, size_t size) { + // CHECK: return hiprandGenerateNormal(gen, data, size, 0.0f, 1.0f); + [](curandGenerator_t gen, float * data, size_t size) { + return curandGenerateNormal(gen, data, size, 0.0f, 1.0f); + } + ); + } + if (distribution == "normal-double") + { + run_benchmark(parser, rng_type, + // CHECK: [](hiprandGenerator_t gen, double * data, size_t size) { + // CHECK: return hiprandGenerateNormalDouble(gen, data, size, 0.0, 1.0); + [](curandGenerator_t gen, double * data, size_t size) { + return curandGenerateNormalDouble(gen, data, size, 0.0, 1.0); + } + ); + } + if (distribution == "log-normal-float") + { + run_benchmark(parser, rng_type, + // CHECK: [](hiprandGenerator_t gen, float * data, size_t size) { + // CHECK: return hiprandGenerateLogNormal(gen, data, size, 0.0f, 1.0f); + [](curandGenerator_t gen, float * data, size_t size) { + return curandGenerateLogNormal(gen, data, size, 0.0f, 1.0f); + } + ); + } + if (distribution == "log-normal-double") + { + run_benchmark(parser, rng_type, + // CHECK: [](hiprandGenerator_t gen, double * data, size_t size) { + // CHECK: return hiprandGenerateLogNormalDouble(gen, data, size, 0.0, 1.0); + [](curandGenerator_t gen, double * data, size_t size) { + return curandGenerateLogNormalDouble(gen, data, size, 0.0, 1.0); + } + ); + } + if (distribution == "poisson") + { + const auto lambdas = parser.get>("lambda"); + for (double lambda : lambdas) + { + std::cout << " " << "lambda " + << std::fixed << std::setprecision(1) << lambda << std::endl; + run_benchmark(parser, rng_type, + // CHECK: [lambda](hiprandGenerator_t gen, unsigned int * data, size_t size) { + // CHECK: return hiprandGeneratePoisson(gen, data, size, lambda); + [lambda](curandGenerator_t gen, unsigned int * data, size_t size) { + return curandGeneratePoisson(gen, data, size, lambda); + } + ); + } + } +} + +const std::vector all_engines = { + "xorwow", + "mrg32k3a", + "mtgp32", + // "mt19937", + "philox", + "sobol32", + // "scrambled_sobol32", + // "sobol64", + // "scrambled_sobol64", +}; + +const std::vector all_distributions = { + "uniform-uint", + "uniform-long-long", + "uniform-float", + "uniform-double", + "normal-float", + "normal-double", + "log-normal-float", + "log-normal-double", + "poisson" +}; + +int main(int argc, char *argv[]) +{ + cli::Parser parser(argc, argv); + + const std::string distribution_desc = + "space-separated list of distributions:" + + std::accumulate(all_distributions.begin(), all_distributions.end(), std::string(), + [](std::string a, std::string b) { + return a + "\n " + b; + } + ) + + "\n or all"; + const std::string engine_desc = + "space-separated list of random number engines:" + + std::accumulate(all_engines.begin(), all_engines.end(), std::string(), + [](std::string a, std::string b) { + return a + "\n " + b; + } + ) + + "\n or all"; + + parser.set_optional("size", "size", DEFAULT_RAND_N, "number of values"); + parser.set_optional("dimensions", "dimensions", 1, "number of dimensions of quasi-random values"); + parser.set_optional("trials", "trials", 20, "number of trials"); + parser.set_optional>("dis", "dis", {"uniform-uint"}, distribution_desc.c_str()); + parser.set_optional>("engine", "engine", {"philox"}, engine_desc.c_str()); + parser.set_optional>("lambda", "lambda", {10.0}, "space-separated list of lambdas of Poisson distribution"); + parser.run_and_exit_if_error(); + + std::vector engines; + { + auto es = parser.get>("engine"); + if (std::find(es.begin(), es.end(), "all") != es.end()) + { + engines = all_engines; + } + else + { + for (auto e : all_engines) + { + if (std::find(es.begin(), es.end(), e) != es.end()) + engines.push_back(e); + } + } + } + + std::vector distributions; + { + auto ds = parser.get>("dis"); + if (std::find(ds.begin(), ds.end(), "all") != ds.end()) + { + distributions = all_distributions; + } + else + { + for (auto d : all_distributions) + { + if (std::find(ds.begin(), ds.end(), d) != ds.end()) + distributions.push_back(d); + } + } + } + + int version; + // CHECK: CURAND_CALL(hiprandGetVersion(&version)); + CURAND_CALL(curandGetVersion(&version)); + int runtime_version; + // cudaRuntimeGetVersion is yet unsupported by HIP + // CHECK-NOT: CUDA_CALL(hipRuntimeGetVersion(&runtime_version)); + CUDA_CALL(cudaRuntimeGetVersion(&runtime_version)); + int device_id; + // CHECK: CUDA_CALL(hipGetDevice(&device_id)); + // CHECK: hipDeviceProp_t props; + // CHECK: CUDA_CALL(hipGetDeviceProperties(&props, device_id)); + CUDA_CALL(cudaGetDevice(&device_id)); + cudaDeviceProp props; + CUDA_CALL(cudaGetDeviceProperties(&props, device_id)); + + std::cout << "cuRAND: " << version << " "; + std::cout << "Runtime: " << runtime_version << " "; + std::cout << "Device: " << props.name; + std::cout << std::endl << std::endl; + + for (auto engine : engines) + { + // CHECK: rng_type_t rng_type = HIPRAND_RNG_PSEUDO_XORWOW; + // CHECK: rng_type = HIPRAND_RNG_PSEUDO_XORWOW; + // CHECK: rng_type = HIPRAND_RNG_PSEUDO_MRG32K3A; + // CHECK: rng_type = HIPRAND_RNG_PSEUDO_MTGP32; + // CHECK: rng_type = HIPRAND_RNG_PSEUDO_MT19937; + // CHECK: rng_type = HIPRAND_RNG_PSEUDO_PHILOX4_32_10; + // CHECK: rng_type = HIPRAND_RNG_QUASI_SOBOL32; + // CHECK: rng_type = HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL32; + // CHECK: rng_type = HIPRAND_RNG_QUASI_SOBOL64; + // CHECK: rng_type = HIPRAND_RNG_QUASI_SCRAMBLED_SOBOL64; + rng_type_t rng_type = CURAND_RNG_PSEUDO_XORWOW; + if (engine == "xorwow") + rng_type = CURAND_RNG_PSEUDO_XORWOW; + else if (engine == "mrg32k3a") + rng_type = CURAND_RNG_PSEUDO_MRG32K3A; + else if (engine == "mtgp32") + rng_type = CURAND_RNG_PSEUDO_MTGP32; + else if (engine == "mt19937") + rng_type = CURAND_RNG_PSEUDO_MT19937; + else if (engine == "philox") + rng_type = CURAND_RNG_PSEUDO_PHILOX4_32_10; + else if (engine == "sobol32") + rng_type = CURAND_RNG_QUASI_SOBOL32; + else if (engine == "scrambled_sobol32") + rng_type = CURAND_RNG_QUASI_SCRAMBLED_SOBOL32; + else if (engine == "sobol64") + rng_type = CURAND_RNG_QUASI_SOBOL64; + else if (engine == "scrambled_sobol64") + rng_type = CURAND_RNG_QUASI_SCRAMBLED_SOBOL64; + else + { + std::cout << "Wrong engine name" << std::endl; + exit(1); + } + + std::cout << engine << ":" << std::endl; + + for (auto distribution : distributions) + { + std::cout << " " << distribution << ":" << std::endl; + run_benchmarks(parser, rng_type, distribution); + } + std::cout << std::endl; + } + + return 0; +} diff --git a/tests/hipify-clang/cuRAND/benchmark_curand_kernel.cpp b/tests/hipify-clang/cuRAND/benchmark_curand_kernel.cpp new file mode 100644 index 0000000000..2c77f25331 --- /dev/null +++ b/tests/hipify-clang/cuRAND/benchmark_curand_kernel.cpp @@ -0,0 +1,669 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// Copyright (c) 2017 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "cmdparser.hpp" +// CHECK: #include +#include +// CHECK: #include +#include +// CHECK: #include +#include +// CHECK-NOT: #include +// CHECK-NOT: #include +#include +#include + +// CHECK: hipError_t error = (x); +// CHECK: if(error!=hipSuccess) { +#define CUDA_CALL(x) do { \ + cudaError_t error = (x);\ + if(error!=cudaSuccess) { \ + printf("Error %d at %s:%d\n",error,__FILE__,__LINE__);\ + exit(EXIT_FAILURE);}} while(0) +#define CURAND_CALL(x) do { if((x)!=CURAND_STATUS_SUCCESS) { \ + printf("Error at %s:%d\n",__FILE__,__LINE__);\ + exit(EXIT_FAILURE);}} while(0) + +#ifndef DEFAULT_RAND_N +const size_t DEFAULT_RAND_N = 1024 * 1024 * 128; +#endif + +size_t next_power2(size_t x) +{ + size_t power = 1; + while (power < x) + { + power *= 2; + } + return power; +} + +template +__global__ +void init_kernel(GeneratorState * states, + const unsigned long long seed, + const unsigned long long offset) +{ + const unsigned int state_id = blockIdx.x * blockDim.x + threadIdx.x; + GeneratorState state; + // CHECK: hiprand_init(seed, state_id, offset, &state); + curand_init(seed, state_id, offset, &state); + states[state_id] = state; +} + +template +__global__ +void generate_kernel(GeneratorState * states, + T * data, + const size_t size, + const GenerateFunc& generate_func, + const Extra extra) +{ + const unsigned int state_id = blockIdx.x * blockDim.x + threadIdx.x; + const unsigned int stride = gridDim.x * blockDim.x; + + GeneratorState state = states[state_id]; + unsigned int index = state_id; + while(index < size) + { + data[index] = generate_func(&state, extra); + index += stride; + } + states[state_id] = state; +} + +template +struct runner +{ + GeneratorState * states; + + runner(const size_t dimensions, + const size_t blocks, + const size_t threads, + const unsigned long long seed, + const unsigned long long offset) + { + const size_t states_size = blocks * threads; + // CHECK: CUDA_CALL(hipMalloc((void **)&states, states_size * sizeof(GeneratorState))); + CUDA_CALL(cudaMalloc((void **)&states, states_size * sizeof(GeneratorState))); + // CHECK: hipLaunchKernelGGL(init_kernel, dim3(blocks), dim3(threads), 0, 0, states, seed, offset); + init_kernel<<>>(states, seed, offset); + // CHECK: CUDA_CALL(hipPeekAtLastError()); + // CHECK: CUDA_CALL(hipDeviceSynchronize()); + CUDA_CALL(cudaPeekAtLastError()); + CUDA_CALL(cudaDeviceSynchronize()); + } + + ~runner() + { + CUDA_CALL(cudaFree(states)); + } + + template + void generate(const size_t blocks, + const size_t threads, + T * data, + const size_t size, + const GenerateFunc& generate_func, + const Extra extra) + { + // CHECK: hipLaunchKernelGGL(generate_kernel, dim3(blocks), dim3(threads), 0, 0, states, data, size, generate_func, extra); + generate_kernel<<>>(states, data, size, generate_func, extra); + } +}; + +// CHECK: void generate_kernel(hiprandStateMtgp32_t * states, +template +__global__ +void generate_kernel(curandStateMtgp32_t * states, + T * data, + const size_t size, + const GenerateFunc& generate_func, + const Extra extra) +{ + const unsigned int state_id = blockIdx.x; + const unsigned int thread_id = threadIdx.x; + unsigned int index = blockIdx.x * blockDim.x + threadIdx.x; + unsigned int stride = gridDim.x * blockDim.x; + // CHECK: __shared__ hiprandStateMtgp32_t state; + __shared__ curandStateMtgp32_t state; + + if (thread_id == 0) + state = states[state_id]; + __syncthreads(); + + const size_t r = size%blockDim.x; + const size_t size_rounded_up = r == 0 ? size : size + (blockDim.x - r); + while(index < size_rounded_up) + { + auto value = generate_func(&state, extra); + if(index < size) + data[index] = value; + index += stride; + } + __syncthreads(); + + if (thread_id == 0) + states[state_id] = state; +} + +// CHECK: struct runner +template<> +struct runner +{ + // CHECK: hiprandStateMtgp32_t * states; + curandStateMtgp32_t * states; + mtgp32_kernel_params_t * d_param; + + runner(const size_t dimensions, + const size_t blocks, + const size_t threads, + const unsigned long long seed, + const unsigned long long offset) + { + const size_t states_size = std::min((size_t)200, blocks); + // CHECK: CUDA_CALL(hipMalloc((void **)&states, states_size * sizeof(hiprandStateMtgp32_t))); + CUDA_CALL(cudaMalloc((void **)&states, states_size * sizeof(curandStateMtgp32_t))); + // CHECK: CUDA_CALL(hipMalloc((void **)&d_param, sizeof(mtgp32_kernel_params))); + CUDA_CALL(cudaMalloc((void **)&d_param, sizeof(mtgp32_kernel_params))); + // curandMakeMTGP32Constants is yet unsupported by HIP + // CHECK-NOT: CURAND_CALL(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, d_param)); + CURAND_CALL(curandMakeMTGP32Constants(mtgp32dc_params_fast_11213, d_param)); + // curandMakeMTGP32KernelState is yet unsupported by HIP + // CHECK-NOT: CURAND_CALL(hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param, states_size, seed)); + CURAND_CALL(curandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param, states_size, seed)); + } + + ~runner() + { + // CHECK: CUDA_CALL(hipFree(states)); + // CHECK: CUDA_CALL(hipFree(d_param)); + CUDA_CALL(cudaFree(states)); + CUDA_CALL(cudaFree(d_param)); + } + + template + void generate(const size_t blocks, + const size_t threads, + T * data, + const size_t size, + const GenerateFunc& generate_func, + const Extra extra) + { + // CHECK: hipLaunchKernelGGL(generate_kernel, dim3(std::min((size_t)200, blocks)), dim3(256), 0, 0, states, data, size, generate_func, extra); + generate_kernel<<>>(states, data, size, generate_func, extra); + } +}; + +// CHECK: void init_kernel(hiprandStateSobol32_t * states, +template +__global__ +void init_kernel(curandStateSobol32_t * states, + const Directions directions, + const unsigned long long offset) +{ + const unsigned int dimension = blockIdx.y; + const unsigned int state_id = blockIdx.x * blockDim.x + threadIdx.x; + // CHECK: hiprandStateSobol32_t state; + // CHECK: hiprand_init(directions[dimension], offset + state_id, &state); + curandStateSobol32_t state; + curand_init(directions[dimension], offset + state_id, &state); + states[gridDim.x * blockDim.x * dimension + state_id] = state; +} + +// CHECK: void generate_kernel(hiprandStateSobol32_t * states, +template +__global__ +void generate_kernel(curandStateSobol32_t * states, + T * data, + const size_t size, + const GenerateFunc& generate_func, + const Extra extra) +{ + const unsigned int dimension = blockIdx.y; + const unsigned int state_id = blockIdx.x * blockDim.x + threadIdx.x; + const unsigned int stride = gridDim.x * blockDim.x; + // CHECK: hiprandStateSobol32_t state = states[gridDim.x * blockDim.x * dimension + state_id]; + curandStateSobol32_t state = states[gridDim.x * blockDim.x * dimension + state_id]; + const unsigned int offset = dimension * size; + unsigned int index = state_id; + while(index < size) + { + data[offset + index] = generate_func(&state, extra); + skipahead(stride - 1, &state); + index += stride; + } + state = states[gridDim.x * blockDim.x * dimension + state_id]; + skipahead(static_cast(size), &state); + states[gridDim.x * blockDim.x * dimension + state_id] = state; +} + +// CHECK: struct runner +template<> +struct runner +{ + // CHECK: hiprandStateSobol32_t * states; + curandStateSobol32_t * states; + size_t dimensions; + + runner(const size_t dimensions, + const size_t blocks, + const size_t threads, + const unsigned long long seed, + const unsigned long long offset) + { + this->dimensions = dimensions; + // CHECK: CUDA_CALL(hipMalloc((void **)&states, states_size * sizeof(hiprandStateSobol32_t))); + const size_t states_size = blocks * threads * dimensions; + CUDA_CALL(cudaMalloc((void **)&states, states_size * sizeof(curandStateSobol32_t))); + // CHECK: hiprandDirectionVectors32_t * directions; + curandDirectionVectors32_t * directions; + // CHECK: const size_t size = dimensions * sizeof(hiprandDirectionVectors32_t); + const size_t size = dimensions * sizeof(curandDirectionVectors32_t); + // CHECK: CUDA_CALL(hipMalloc((void **)&directions, size)); + 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)); + 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)); + + const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions); + // CHECK: hipLaunchKernelGGL(init_kernel, dim3(dim3(blocks_x, dimensions)), dim3(threads), 0, 0, states, directions, offset); + init_kernel<<>>(states, directions, offset); + // CHECK: CUDA_CALL(hipPeekAtLastError()); + // CHECK: CUDA_CALL(hipDeviceSynchronize()); + CUDA_CALL(cudaPeekAtLastError()); + CUDA_CALL(cudaDeviceSynchronize()); + // CHECK: CUDA_CALL(hipFree(directions)); + CUDA_CALL(cudaFree(directions)); + } + + ~runner() + { + // CHECK: CUDA_CALL(hipFree(states)); + CUDA_CALL(cudaFree(states)); + } + + template + void generate(const size_t blocks, + const size_t threads, + T * data, + const size_t size, + const GenerateFunc& generate_func, + const Extra extra) + { + const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions); + // CHECK: hipLaunchKernelGGL(generate_kernel, dim3(dim3(blocks_x, dimensions)), dim3(threads), 0, 0, states, data, size / dimensions, generate_func, extra); + generate_kernel<<>>(states, data, size / dimensions, generate_func, extra); + } +}; + +template +void run_benchmark(const cli::Parser& parser, + const GenerateFunc& generate_func, + const Extra extra) +{ + const size_t size = parser.get("size"); + const size_t dimensions = parser.get("dimensions"); + const size_t trials = parser.get("trials"); + + const size_t blocks = parser.get("blocks"); + const size_t threads = parser.get("threads"); + + T * data; + // CHECK: CUDA_CALL(hipMalloc((void **)&data, size * sizeof(T))); + CUDA_CALL(cudaMalloc((void **)&data, size * sizeof(T))); + + runner r(dimensions, blocks, threads, 12345ULL, 6789ULL); + + // Warm-up + for (size_t i = 0; i < 5; i++) + { + r.generate(blocks, threads, data, size, generate_func, extra); + // CHECK: CUDA_CALL(hipPeekAtLastError()); + // CHECK: CUDA_CALL(hipDeviceSynchronize()); + CUDA_CALL(cudaPeekAtLastError()); + CUDA_CALL(cudaDeviceSynchronize()); + } + // CHECK: CUDA_CALL(hipDeviceSynchronize()); + CUDA_CALL(cudaDeviceSynchronize()); + + // Measurement + auto start = std::chrono::high_resolution_clock::now(); + for (size_t i = 0; i < trials; i++) + { + r.generate(blocks, threads, data, size, generate_func, extra); + } + // CHECK: CUDA_CALL(hipPeekAtLastError()); + // CHECK: CUDA_CALL(hipDeviceSynchronize()); + CUDA_CALL(cudaPeekAtLastError()); + CUDA_CALL(cudaDeviceSynchronize()); + auto end = std::chrono::high_resolution_clock::now(); + std::chrono::duration elapsed = end - start; + + std::cout << std::fixed << std::setprecision(3) + << " " + << "Throughput = " + << std::setw(8) << (trials * size * sizeof(T)) / + (elapsed.count() / 1e3 * (1 << 30)) + << " GB/s, Samples = " + << std::setw(8) << (trials * size) / + (elapsed.count() / 1e3 * (1 << 30)) + << " GSample/s, AvgTime (1 trial) = " + << std::setw(8) << elapsed.count() / trials + << " ms, Time (all) = " + << std::setw(8) << elapsed.count() + << " ms, Size = " << size + << std::endl; + // CHECK: CUDA_CALL(hipFree(data)); + CUDA_CALL(cudaFree(data)); +} + +template +void run_benchmarks(const cli::Parser& parser, + const std::string& distribution) +{ + if (distribution == "uniform-uint") + { + // curandStateSobol64_t and curandStateScrambledSobol64_t are yet unsupported by HIP + // CHECK-NOT: if (!std::is_same::value && + // CHECK-NOT: !std::is_same::value) + if (!std::is_same::value && + !std::is_same::value) + { + run_benchmark(parser, + [] __device__ (GeneratorState * state, int) { + // CHECK: return hiprand(state); + return curand(state); + }, 0 + ); + } + } + if (distribution == "uniform-long-long") + { + // curandStateSobol64_t and curandStateScrambledSobol64_t are yet unsupported by HIP + // CHECK-NOT: if (!std::is_same::value && + // CHECK-NOT: !std::is_same::value) + if (std::is_same::value || + std::is_same::value) + { + run_benchmark(parser, + [] __device__ (GeneratorState * state, int) { + // CHECK: return hiprand(state); + return curand(state); + }, 0 + ); + } + } + if (distribution == "uniform-float") + { + run_benchmark(parser, + [] __device__ (GeneratorState * state, int) { + // CHECK: return hiprand_uniform(state); + return curand_uniform(state); + }, 0 + ); + } + if (distribution == "uniform-double") + { + run_benchmark(parser, + [] __device__ (GeneratorState * state, int) { + // CHECK: return hiprand_uniform_double(state); + return curand_uniform_double(state); + }, 0 + ); + } + if (distribution == "normal-float") + { + run_benchmark(parser, + [] __device__ (GeneratorState * state, int) { + // CHECK: return hiprand_normal(state); + return curand_normal(state); + }, 0 + ); + } + if (distribution == "normal-double") + { + run_benchmark(parser, + [] __device__ (GeneratorState * state, int) { + // CHECK: return hiprand_normal_double(state); + return curand_normal_double(state); + }, 0 + ); + } + if (distribution == "log-normal-float") + { + run_benchmark(parser, + [] __device__ (GeneratorState * state, int) { + // CHECK: return hiprand_log_normal(state, 0.0f, 1.0f); + return curand_log_normal(state, 0.0f, 1.0f); + }, 0 + ); + } + if (distribution == "log-normal-double") + { + run_benchmark(parser, + [] __device__ (GeneratorState * state, int) { + // CHECK: return hiprand_log_normal_double(state, 0.0, 1.0); + return curand_log_normal_double(state, 0.0, 1.0); + }, 0 + ); + } + if (distribution == "poisson") + { + const auto lambdas = parser.get>("lambda"); + for (double lambda : lambdas) + { + std::cout << " " << "lambda " + << std::fixed << std::setprecision(1) << lambda << std::endl; + run_benchmark(parser, + [] __device__ (GeneratorState * state, double lambda) { + // CHECK: return hiprand_poisson(state, lambda); + return curand_poisson(state, lambda); + }, lambda + ); + } + } + if (distribution == "discrete-poisson") + { + const auto lambdas = parser.get>("lambda"); + for (double lambda : lambdas) + { + std::cout << " " << "lambda " + << std::fixed << std::setprecision(1) << lambda << std::endl; + // CHECK: hiprandDiscreteDistribution_t discrete_distribution; + curandDiscreteDistribution_t discrete_distribution; + // CHECK: CURAND_CALL(hiprandCreatePoissonDistribution(lambda, &discrete_distribution)); + CURAND_CALL(curandCreatePoissonDistribution(lambda, &discrete_distribution)); + run_benchmark(parser, + // CHECK: [] __device__ (GeneratorState * state, hiprandDiscreteDistribution_t discrete_distribution) { + [] __device__ (GeneratorState * state, curandDiscreteDistribution_t discrete_distribution) { + // CHECK: return hiprand_discrete4(state, discrete_distribution); + return curand_discrete(state, discrete_distribution); + }, discrete_distribution + ); + // CHECK: CURAND_CALL(hiprandDestroyDistribution(discrete_distribution)); + CURAND_CALL(curandDestroyDistribution(discrete_distribution)); + } + } +} + +const std::vector all_engines = { + "xorwow", + "mrg32k3a", + "mtgp32", + // "mt19937", + "philox", + "sobol32", + // "scrambled_sobol32", + // "sobol64", + // "scrambled_sobol64", +}; + +const std::vector all_distributions = { + "uniform-uint", + // "uniform-long-long", + "uniform-float", + "uniform-double", + "normal-float", + "normal-double", + "log-normal-float", + "log-normal-double", + "poisson", + "discrete-poisson", +}; + +int main(int argc, char *argv[]) +{ + cli::Parser parser(argc, argv); + + const std::string distribution_desc = + "space-separated list of distributions:" + + std::accumulate(all_distributions.begin(), all_distributions.end(), std::string(), + [](std::string a, std::string b) { + return a + "\n " + b; + } + ) + + "\n or all"; + const std::string engine_desc = + "space-separated list of random number engines:" + + std::accumulate(all_engines.begin(), all_engines.end(), std::string(), + [](std::string a, std::string b) { + return a + "\n " + b; + } + ) + + "\n or all"; + + parser.set_optional("size", "size", DEFAULT_RAND_N, "number of values"); + parser.set_optional("dimensions", "dimensions", 1, "number of dimensions of quasi-random values"); + parser.set_optional("trials", "trials", 20, "number of trials"); + parser.set_optional("blocks", "blocks", 256, "number of blocks"); + parser.set_optional("threads", "threads", 256, "number of threads in each block"); + parser.set_optional>("dis", "dis", {"uniform-uint"}, distribution_desc.c_str()); + parser.set_optional>("engine", "engine", {"philox"}, engine_desc.c_str()); + parser.set_optional>("lambda", "lambda", {10.0}, "space-separated list of lambdas of Poisson distribution"); + parser.run_and_exit_if_error(); + + std::vector engines; + { + auto es = parser.get>("engine"); + if (std::find(es.begin(), es.end(), "all") != es.end()) + { + engines = all_engines; + } + else + { + for (auto e : all_engines) + { + if (std::find(es.begin(), es.end(), e) != es.end()) + engines.push_back(e); + } + } + } + + std::vector distributions; + { + auto ds = parser.get>("dis"); + if (std::find(ds.begin(), ds.end(), "all") != ds.end()) + { + distributions = all_distributions; + } + else + { + for (auto d : all_distributions) + { + if (std::find(ds.begin(), ds.end(), d) != ds.end()) + distributions.push_back(d); + } + } + } + + int version; + // CHECK: CURAND_CALL(hiprandGetVersion(&version)); + CURAND_CALL(curandGetVersion(&version)); + int runtime_version; + // cudaRuntimeGetVersion is yet unsupported by HIP + // CHECK-NOT: CUDA_CALL(hipRuntimeGetVersion(&runtime_version)); + CUDA_CALL(cudaRuntimeGetVersion(&runtime_version)); + int device_id; + // CHECK: CUDA_CALL(hipGetDevice(&device_id)); + // CHECK: hipDeviceProp_t props; + // CHECK: CUDA_CALL(hipGetDeviceProperties(&props, device_id)); + CUDA_CALL(cudaGetDevice(&device_id)); + cudaDeviceProp props; + CUDA_CALL(cudaGetDeviceProperties(&props, device_id)); + + std::cout << "cuRAND: " << version << " "; + std::cout << "Runtime: " << runtime_version << " "; + std::cout << "Device: " << props.name; + std::cout << std::endl << std::endl; + + for (auto engine : engines) + { + std::cout << engine << ":" << std::endl; + for (auto distribution : distributions) + { + std::cout << " " << distribution << ":" << std::endl; + const std::string plot_name = engine + "-" + distribution; + if (engine == "xorwow") + { + // CHECK: run_benchmarks(parser, distribution); + run_benchmarks(parser, distribution); + } + else if (engine == "mrg32k3a") + { + // CHECK: run_benchmarks(parser, distribution); + run_benchmarks(parser, distribution); + } + else if (engine == "philox") + { + // CHECK: run_benchmarks(parser, distribution); + run_benchmarks(parser, distribution); + } + else if (engine == "sobol32") + { + // CHECK: run_benchmarks(parser, distribution); + run_benchmarks(parser, distribution); + } + else if (engine == "mtgp32") + { + // CHECK: run_benchmarks(parser, distribution); + run_benchmarks(parser, distribution); + } + } + } + + return 0; +} diff --git a/tests/hipify-clang/cuRAND/cmdparser.hpp b/tests/hipify-clang/cuRAND/cmdparser.hpp new file mode 100644 index 0000000000..364612ebb1 --- /dev/null +++ b/tests/hipify-clang/cuRAND/cmdparser.hpp @@ -0,0 +1,513 @@ +// The MIT License (MIT) +// +// Copyright (c) 2015 - 2016 Florian Rappl +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +/* + This file is part of the C++ CmdParser utility. + Copyright (c) 2015 - 2016 Florian Rappl +*/ + +#pragma once +#include +#include +#include +#include +#include +#include + +namespace cli { + struct CallbackArgs { + const std::vector& arguments; + std::ostream& output; + std::ostream& error; + }; + class Parser { + private: + class CmdBase { + public: + explicit CmdBase(const std::string& name, const std::string& alternative, const std::string& description, bool required, bool dominant, bool variadic) : + name(name), + command(name.size() > 0 ? "-" + name : ""), + alternative(alternative.size() > 0 ? "--" + alternative : ""), + description(description), + required(required), + handled(false), + arguments({}), + dominant(dominant), + variadic(variadic) { + } + + virtual ~CmdBase() { + } + + std::string name; + std::string command; + std::string alternative; + std::string description; + bool required; + bool handled; + std::vector arguments; + bool const dominant; + bool const variadic; + + virtual std::string print_value() const = 0; + virtual bool parse(std::ostream& output, std::ostream& error) = 0; + + bool is(const std::string& given) const { + return given == command || given == alternative; + } + }; + + template + struct ArgumentCountChecker + { + static constexpr bool Variadic = false; + }; + + template + struct ArgumentCountChecker> + { + static constexpr bool Variadic = true; + }; + + template + class CmdFunction final : public CmdBase { + public: + explicit CmdFunction(const std::string& name, const std::string& alternative, const std::string& description, bool required, bool dominant) : + CmdBase(name, alternative, description, required, dominant, ArgumentCountChecker::Variadic) { + } + + virtual bool parse(std::ostream& output, std::ostream& error) { + try { + CallbackArgs args { arguments, output, error }; + value = callback(args); + return true; + } catch (...) { + return false; + } + } + + virtual std::string print_value() const { + return ""; + } + + std::function callback; + T value; + }; + + template + class CmdArgument final : public CmdBase { + public: + explicit CmdArgument(const std::string& name, const std::string& alternative, const std::string& description, bool required, bool dominant) : + CmdBase(name, alternative, description, required, dominant, ArgumentCountChecker::Variadic) { + } + + virtual bool parse(std::ostream&, std::ostream&) { + try { + value = Parser::parse(arguments, value); + return true; + } catch (...) { + return false; + } + } + + virtual std::string print_value() const { + return stringify(value); + } + + T value; + }; + + static int parse(const std::vector& elements, const int&) { + if (elements.size() != 1) + throw std::bad_cast(); + + return std::stoi(elements[0]); + } + + static bool parse(const std::vector& elements, const bool& defval) { + if (elements.size() != 0) + throw std::runtime_error("A boolean command line parameter cannot have any arguments."); + + return !defval; + } + + static double parse(const std::vector& elements, const double&) { + if (elements.size() != 1) + throw std::bad_cast(); + + return std::stod(elements[0]); + } + + static float parse(const std::vector& elements, const float&) { + if (elements.size() != 1) + throw std::bad_cast(); + + return std::stof(elements[0]); + } + + static long double parse(const std::vector& elements, const long double&) { + if (elements.size() != 1) + throw std::bad_cast(); + + return std::stold(elements[0]); + } + + static unsigned int parse(const std::vector& elements, const unsigned int&) { + if (elements.size() != 1) + throw std::bad_cast(); + + return static_cast(std::stoul(elements[0])); + } + + static unsigned long parse(const std::vector& elements, const unsigned long&) { + if (elements.size() != 1) + throw std::bad_cast(); + + return std::stoul(elements[0]); + } + + static unsigned long long parse(const std::vector& elements, const unsigned long long&) { + if (elements.size() != 1) + throw std::bad_cast(); + + return std::stoull(elements[0]); + } + + static long parse(const std::vector& elements, const long&) { + if (elements.size() != 1) + throw std::bad_cast(); + + return std::stol(elements[0]); + } + + static std::string parse(const std::vector& elements, const std::string&) { + if (elements.size() != 1) + throw std::bad_cast(); + + return elements[0]; + } + + template + static std::vector parse(const std::vector& elements, const std::vector&) { + const T defval = T(); + std::vector values { }; + std::vector buffer(1); + + for (const auto& element : elements) { + buffer[0] = element; + values.push_back(parse(buffer, defval)); + } + + return values; + } + + template + static std::string stringify(const T& value) { + return std::to_string(value); + } + + template + static std::string stringify(const std::vector& values) { + std::stringstream ss { }; + ss << "[ "; + + for (const auto& value : values) { + ss << stringify(value) << " "; + } + + ss << "]"; + return ss.str(); + } + + static std::string stringify(const std::string& str) { + return str; + } + + public: + explicit Parser(int argc, const char** argv) : + _appname(argv[0]) { + for (int i = 1; i < argc; ++i) { + _arguments.push_back(argv[i]); + } + enable_help(); + } + + explicit Parser(int argc, char** argv) : + _appname(argv[0]) { + for (int i = 1; i < argc; ++i) { + _arguments.push_back(argv[i]); + } + enable_help(); + } + + ~Parser() { + for (int i = 0, n = _commands.size(); i < n; ++i) { + delete _commands[i]; + } + } + + bool has_help() const { + for (const auto command : _commands) { + if (command->name == "h" && command->alternative == "--help") { + return true; + } + } + + return false; + } + + void enable_help() { + set_callback("h", "help", std::function([this](CallbackArgs& args){ + args.output << this->usage(); + exit(0); + return false; + }), "", true); + } + + void disable_help() { + for (auto command = _commands.begin(); command != _commands.end(); ++command) { + if ((*command)->name == "h" && (*command)->alternative == "--help") { + _commands.erase(command); + break; + } + } + } + + template + void set_default(bool is_required, const std::string& description = "") { + auto command = new CmdArgument { "", "", description, is_required, false }; + _commands.push_back(command); + } + + template + void set_required(const std::string& name, const std::string& alternative, const std::string& description = "", bool dominant = false) { + auto command = new CmdArgument { name, alternative, description, true, dominant }; + _commands.push_back(command); + } + + template + void set_optional(const std::string& name, const std::string& alternative, T defaultValue, const std::string& description = "", bool dominant = false) { + auto command = new CmdArgument { name, alternative, description, false, dominant }; + command->value = defaultValue; + _commands.push_back(command); + } + + template + void set_callback(const std::string& name, const std::string& alternative, std::function callback, const std::string& description = "", bool dominant = false) { + auto command = new CmdFunction { name, alternative, description, false, dominant }; + command->callback = callback; + _commands.push_back(command); + } + + inline void run_and_exit_if_error() { + if (run() == false) { + exit(1); + } + } + + inline bool run() { + return run(std::cout, std::cerr); + } + + inline bool run(std::ostream& output) { + return run(output, std::cerr); + } + + bool run(std::ostream& output, std::ostream& error) { + if (_arguments.size() > 0) { + auto current = find_default(); + + for (int i = 0, n = _arguments.size(); i < n; ++i) { + auto isarg = _arguments[i].size() > 0 && _arguments[i][0] == '-'; + auto associated = isarg ? find(_arguments[i]) : nullptr; + + if (associated != nullptr) { + current = associated; + associated->handled = true; + } else if (current == nullptr) { + error << no_default(); + return false; + } else { + current->arguments.push_back(_arguments[i]); + current->handled = true; + if (!current->variadic) + { + // If the current command is not variadic, then no more arguments + // should be added to it. In this case, switch back to the default + // command. + current = find_default(); + } + } + } + } + + // First, parse dominant arguments since they succeed even if required + // arguments are missing. + for (auto command : _commands) { + if (command->handled && command->dominant && !command->parse(output, error)) { + error << howto_use(command); + return false; + } + } + + // Next, check for any missing arguments. + for (auto command : _commands) { + if (command->required && !command->handled) { + error << howto_required(command); + return false; + } + } + + // Finally, parse all remaining arguments. + for (auto command : _commands) { + if (command->handled && !command->dominant && !command->parse(output, error)) { + error << howto_use(command); + return false; + } + } + + return true; + } + + template + T get(const std::string& name) const { + for (const auto& command : _commands) { + if (command->name == name) { + auto cmd = dynamic_cast*>(command); + + if (cmd == nullptr) { + throw std::runtime_error("Invalid usage of the parameter " + name + " detected."); + } + + return cmd->value; + } + } + + throw std::runtime_error("The parameter " + name + " could not be found."); + } + + template + T get_if(const std::string& name, std::function callback) const { + auto value = get(name); + return callback(value); + } + + int requirements() const { + int count = 0; + + for (const auto& command : _commands) { + if (command->required) { + ++count; + } + } + + return count; + } + + int commands() const { + return static_cast(_commands.size()); + } + + inline const std::string& app_name() const { + return _appname; + } + + protected: + CmdBase* find(const std::string& name) { + for (auto command : _commands) { + if (command->is(name)) { + return command; + } + } + + return nullptr; + } + + CmdBase* find_default() { + for (auto command : _commands) { + if (command->name == "") { + return command; + } + } + + return nullptr; + } + + std::string usage() const { + std::stringstream ss { }; + ss << "Available parameters:\n\n"; + + for (const auto& command : _commands) { + ss << " " << command->command << "\t" << command->alternative; + + if (command->required == true) { + ss << "\t(required)"; + } + + ss << "\n " << command->description; + + if (command->required == false) { + ss << "\n " << "This parameter is optional. The default value is '" + command->print_value() << "'."; + } + + ss << "\n\n"; + } + + return ss.str(); + } + + void print_help(std::stringstream& ss) const { + if (has_help()) { + ss << "For more help use --help or -h.\n"; + } + } + + std::string howto_required(CmdBase* command) const { + std::stringstream ss { }; + ss << "The parameter " << command->name << " is required.\n"; + ss << command->description << '\n'; + print_help(ss); + return ss.str(); + } + + std::string howto_use(CmdBase* command) const { + std::stringstream ss { }; + ss << "The parameter " << command->name << " has invalid arguments.\n"; + ss << command->description << '\n'; + print_help(ss); + return ss.str(); + } + + std::string no_default() const { + std::stringstream ss { }; + ss << "No default parameter has been specified.\n"; + ss << "The given argument must be used with a parameter.\n"; + print_help(ss); + return ss.str(); + } + + private: + const std::string _appname; + std::vector _arguments; + std::vector _commands; + }; +} diff --git a/tests/hipify-clang/headers_test_09.cu b/tests/hipify-clang/headers_test_09.cu new file mode 100644 index 0000000000..048ac2e9a9 --- /dev/null +++ b/tests/hipify-clang/headers_test_09.cu @@ -0,0 +1,93 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #include +// CHECK: #include + +// CHECK-NOT: #include +// CHECK-NOT: #include + +// CHECK: #include "hip/hip_runtime_api.h" +// CHECK: #include "hip/channel_descriptor.h" +// CHECK: #include "hip/device_functions.h" +// CHECK: #include "hip/driver_types.h" +// CHECK: #include "hip/hip_complex.h" +// CHECK: #include "hip/hip_fp16.h" +// CHECK: #include "hip/hip_texture_types.h" +// CHECK: #include "hip/hip_vector_types.h" + +// CHECK: #include + +// CHECK: #include "hipblas.h" +// CHECK-NOT: #include "cublas.h" + +// CHECK: #include + +// CHECK: #include "hiprand.h" +// CHECK: #include "hiprand_kernel.h" + +// CHECK: #include + +// CHECK-NOT: #include "hiprand.h" +// CHECK-NOT: #include "hiprand_kernel.h" +// CHECK-NOT: #include "curand_discrete.h" +// CHECK-NOT: #include "curand_discrete2.h" +// CHECK-NOT: #include "curand_globals.h" +// CHECK-NOT: #include "curand_lognormal.h" +// CHECK-NOT: #include "curand_mrg32k3a.h" +// CHECK-NOT: #include "curand_mtgp32.h" +// CHECK-NOT: #include "curand_mtgp32_host.h" +// CHECK-NOT: #include "curand_mtgp32_kernel.h" +// CHECK-NOT: #include "curand_mtgp32dc_p_11213.h" +// CHECK-NOT: #include "curand_normal.h" +// CHECK-NOT: #include "curand_normal_static.h" +// CHECK-NOT: #include "curand_philox4x32_x.h" +// CHECK-NOT: #include "curand_poisson.h" +// CHECK-NOT: #include "curand_precalc.h" +// CHECK-NOT: #include "curand_uniform.h" + +// CHECK: #include + +#include + +#include + +#include + +#include "cuda_runtime_api.h" +#include "channel_descriptor.h" +#include "device_functions.h" +#include "driver_types.h" +#include "cuComplex.h" +#include "cuda_fp16.h" +#include "cuda_texture_types.h" +#include "vector_types.h" + +#include + +#include "cublas_v2.h" +#include "cublas.h" + +#include + +#include "curand.h" +#include "curand_kernel.h" + +#include + +#include "curand_discrete.h" +#include "curand_discrete2.h" +#include "curand_globals.h" +#include "curand_lognormal.h" +#include "curand_mrg32k3a.h" +#include "curand_mtgp32.h" +#include "curand_mtgp32_host.h" +#include "curand_mtgp32_kernel.h" +#include "curand_mtgp32dc_p_11213.h" +#include "curand_normal.h" +#include "curand_normal_static.h" +#include "curand_philox4x32_x.h" +#include "curand_poisson.h" +#include "curand_precalc.h" +#include "curand_uniform.h" + +#include diff --git a/tests/hipify-clang/lit.cfg b/tests/hipify-clang/lit.cfg index 12b1410cee..104cee8311 100644 --- a/tests/hipify-clang/lit.cfg +++ b/tests/hipify-clang/lit.cfg @@ -21,6 +21,8 @@ config.test_format = lit.formats.ShTest() # test_source_root: The root path where tests are located. config.test_source_root = os.path.dirname(__file__) +config.excludes = ['cmdparser.hpp'] + # test_exec_root: The path where tests are located (default is the test suite root). #config.test_exec_root = config.test_source_root