AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

[HIPIFY] InclusionDirective refactoring

Просмотр исходного кода 
Due to support of cuRAND headers.

+ compound test on all headers is added;
+ missing entities are added with updating the doc;
+ a couple cuRAND tests are added (https://github.com/ROCmSoftwarePlatform/rocRAND/tree/master/benchmark):
  - the following CUDA entities are still unsupported by hipRAND:
      curandMakeMTGP32Constants
      curandMakeMTGP32KernelState
      curandGetDirectionVectors32
      curandDirectionVectorSet_t
      CURAND_DIRECTION_VECTORS_32_JOEKUO6
      curandStateSobol64_t
      curandStateScrambledSobol64_t
      curandGenerateLongLong
  - and the following - by HIP:
      cudaRuntimeGetVersion
  - those entities are handled by CHECK-NOT directive for now.


[ROCm/hip commit: 02e23c4d87]
Этот коммит содержится в:
Evgeny Mankov
2018-01-29 18:33:47 +03:00
родитель fad10eed7d
Коммит e455192444
11 изменённых файлов: 1756 добавлений и 42 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/hip/docs/markdown/CURAND_API_supported_by_HIP.md
+2
Просмотреть файл
@@ -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` |
projects/hip/hipify-clang/src/CUDA2HipMap.cpp
+31 -14
Просмотреть файл
@@ -367,24 +367,39 @@ const std::map<llvm::StringRef, hipCounter> CUDA_TYPE_NAME_MAP{
/// Maps cuda header names to hip header names.
const std::map <llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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}},
projects/hip/hipify-clang/src/HipifyAction.cpp
+44 -24
Просмотреть файл
@@ -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);
projects/hip/hipify-clang/src/HipifyAction.h
+5
Просмотреть файл
@@ -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<clang::ASTConsumer> CreateASTConsumer(clang::CompilerInstance &CI, llvm::StringRef InFile) override;
bool Exclude(const hipCounter & hipToken);
};
projects/hip/hipify-clang/src/Statistics.cpp
+2 -2
Просмотреть файл
@@ -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) {
projects/hip/hipify-clang/src/Statistics.h
+2 -2
Просмотреть файл
@@ -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;
};
projects/hip/tests/hipify-clang/cuRAND/benchmark_curand_generate.cpp
+393
Просмотреть файл
@@ -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 <iostream>
#include <iomanip>
#include <vector>
#include <string>
#include <chrono>
#include <numeric>
#include <utility>
#include <algorithm>
#include "cmdparser.hpp"
// CHECK: #include <hip/hip_runtime.h>
#include <cuda_runtime.h>
// CHECK: #include <hiprand.h>
#include <curand.h>
// 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<hiprandStatus_t(hiprandGenerator_t, T *, size_t)>;
template<typename T>
using generate_func_type = std::function<curandStatus_t(curandGenerator_t, T *, size_t)>;
template<typename T>
void run_benchmark(const cli::Parser& parser,
const rng_type_t rng_type,
generate_func_type<T> generate_func)
{
const size_t size = parser.get<size_t>("size");
const size_t trials = parser.get<size_t>("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<size_t>("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<double, std::milli> 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<unsigned int>(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<unsigned long long>(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<float>(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<double>(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<float>(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<double>(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<float>(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<double>(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<std::vector<double>>("lambda");
for (double lambda : lambdas)
{
std::cout << " " << "lambda "
<< std::fixed << std::setprecision(1) << lambda << std::endl;
run_benchmark<unsigned int>(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<std::string> all_engines = {
"xorwow",
"mrg32k3a",
"mtgp32",
// "mt19937",
"philox",
"sobol32",
// "scrambled_sobol32",
// "sobol64",
// "scrambled_sobol64",
};
const std::vector<std::string> 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_t>("size", "size", DEFAULT_RAND_N, "number of values");
parser.set_optional<size_t>("dimensions", "dimensions", 1, "number of dimensions of quasi-random values");
parser.set_optional<size_t>("trials", "trials", 20, "number of trials");
parser.set_optional<std::vector<std::string>>("dis", "dis", {"uniform-uint"}, distribution_desc.c_str());
parser.set_optional<std::vector<std::string>>("engine", "engine", {"philox"}, engine_desc.c_str());
parser.set_optional<std::vector<double>>("lambda", "lambda", {10.0}, "space-separated list of lambdas of Poisson distribution");
parser.run_and_exit_if_error();
std::vector<std::string> engines;
{
auto es = parser.get<std::vector<std::string>>("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<std::string> distributions;
{
auto ds = parser.get<std::vector<std::string>>("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;
}
projects/hip/tests/hipify-clang/cuRAND/benchmark_curand_kernel.cpp
+669
Просмотреть файл
@@ -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 <iostream>
#include <iomanip>
#include <vector>
#include <string>
#include <chrono>
#include <numeric>
#include <utility>
#include <type_traits>
#include <algorithm>
#include "cmdparser.hpp"
// CHECK: #include <hip/hip_runtime.h>
#include <cuda_runtime.h>
// CHECK: #include <hiprand.h>
#include <curand.h>
// CHECK: #include <hiprand_kernel.h>
#include <curand_kernel.h>
// CHECK-NOT: #include <curand_mtgp32_host.h>
// CHECK-NOT: #include <curand_mtgp32dc_p_11213.h>
#include <curand_mtgp32_host.h>
#include <curand_mtgp32dc_p_11213.h>
// 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<typename GeneratorState>
__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<typename GeneratorState, typename T, typename GenerateFunc, typename Extra>
__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<typename GeneratorState>
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<<<blocks, threads>>>(states, seed, offset);
// CHECK: CUDA_CALL(hipPeekAtLastError());
// CHECK: CUDA_CALL(hipDeviceSynchronize());
CUDA_CALL(cudaPeekAtLastError());
CUDA_CALL(cudaDeviceSynchronize());
}
~runner()
{
CUDA_CALL(cudaFree(states));
}
template<typename T, typename GenerateFunc, typename Extra>
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<<<blocks, threads>>>(states, data, size, generate_func, extra);
}
};
// CHECK: void generate_kernel(hiprandStateMtgp32_t * states,
template<typename T, typename GenerateFunc, typename Extra>
__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<hiprandStateMtgp32_t>
template<>
struct runner<curandStateMtgp32_t>
{
// 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<typename T, typename GenerateFunc, typename Extra>
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<<<std::min((size_t)200, blocks), 256>>>(states, data, size, generate_func, extra);
}
};
// CHECK: void init_kernel(hiprandStateSobol32_t * states,
template<typename Directions>
__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<typename T, typename GenerateFunc, typename Extra>
__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<unsigned int>(size), &state);
states[gridDim.x * blockDim.x * dimension + state_id] = state;
}
// CHECK: struct runner<hiprandStateSobol32_t>
template<>
struct runner<curandStateSobol32_t>
{
// 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<<<dim3(blocks_x, dimensions), threads>>>(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<typename T, typename GenerateFunc, typename Extra>
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<<<dim3(blocks_x, dimensions), threads>>>(states, data, size / dimensions, generate_func, extra);
}
};
template<typename T, typename GeneratorState, typename GenerateFunc, typename Extra>
void run_benchmark(const cli::Parser& parser,
const GenerateFunc& generate_func,
const Extra extra)
{
const size_t size = parser.get<size_t>("size");
const size_t dimensions = parser.get<size_t>("dimensions");
const size_t trials = parser.get<size_t>("trials");
const size_t blocks = parser.get<size_t>("blocks");
const size_t threads = parser.get<size_t>("threads");
T * data;
// CHECK: CUDA_CALL(hipMalloc((void **)&data, size * sizeof(T)));
CUDA_CALL(cudaMalloc((void **)&data, size * sizeof(T)));
runner<GeneratorState> 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<double, std::milli> 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<typename GeneratorState>
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<GeneratorState, hiprandStateSobol64_t>::value &&
// CHECK-NOT: !std::is_same<GeneratorState, hiprandStateScrambledSobol64_t>::value)
if (!std::is_same<GeneratorState, curandStateSobol64_t>::value &&
!std::is_same<GeneratorState, curandStateScrambledSobol64_t>::value)
{
run_benchmark<unsigned int, GeneratorState>(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<GeneratorState, hiprandStateSobol64_t>::value &&
// CHECK-NOT: !std::is_same<GeneratorState, hiprandStateScrambledSobol64_t>::value)
if (std::is_same<GeneratorState, curandStateSobol64_t>::value ||
std::is_same<GeneratorState, curandStateScrambledSobol64_t>::value)
{
run_benchmark<unsigned long long, GeneratorState>(parser,
[] __device__ (GeneratorState * state, int) {
// CHECK: return hiprand(state);
return curand(state);
}, 0
);
}
}
if (distribution == "uniform-float")
{
run_benchmark<float, GeneratorState>(parser,
[] __device__ (GeneratorState * state, int) {
// CHECK: return hiprand_uniform(state);
return curand_uniform(state);
}, 0
);
}
if (distribution == "uniform-double")
{
run_benchmark<double, GeneratorState>(parser,
[] __device__ (GeneratorState * state, int) {
// CHECK: return hiprand_uniform_double(state);
return curand_uniform_double(state);
}, 0
);
}
if (distribution == "normal-float")
{
run_benchmark<float, GeneratorState>(parser,
[] __device__ (GeneratorState * state, int) {
// CHECK: return hiprand_normal(state);
return curand_normal(state);
}, 0
);
}
if (distribution == "normal-double")
{
run_benchmark<double, GeneratorState>(parser,
[] __device__ (GeneratorState * state, int) {
// CHECK: return hiprand_normal_double(state);
return curand_normal_double(state);
}, 0
);
}
if (distribution == "log-normal-float")
{
run_benchmark<float, GeneratorState>(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<double, GeneratorState>(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<std::vector<double>>("lambda");
for (double lambda : lambdas)
{
std::cout << " " << "lambda "
<< std::fixed << std::setprecision(1) << lambda << std::endl;
run_benchmark<unsigned int, GeneratorState>(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<std::vector<double>>("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<unsigned int, GeneratorState>(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<std::string> all_engines = {
"xorwow",
"mrg32k3a",
"mtgp32",
// "mt19937",
"philox",
"sobol32",
// "scrambled_sobol32",
// "sobol64",
// "scrambled_sobol64",
};
const std::vector<std::string> 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_t>("size", "size", DEFAULT_RAND_N, "number of values");
parser.set_optional<size_t>("dimensions", "dimensions", 1, "number of dimensions of quasi-random values");
parser.set_optional<size_t>("trials", "trials", 20, "number of trials");
parser.set_optional<size_t>("blocks", "blocks", 256, "number of blocks");
parser.set_optional<size_t>("threads", "threads", 256, "number of threads in each block");
parser.set_optional<std::vector<std::string>>("dis", "dis", {"uniform-uint"}, distribution_desc.c_str());
parser.set_optional<std::vector<std::string>>("engine", "engine", {"philox"}, engine_desc.c_str());
parser.set_optional<std::vector<double>>("lambda", "lambda", {10.0}, "space-separated list of lambdas of Poisson distribution");
parser.run_and_exit_if_error();
std::vector<std::string> engines;
{
auto es = parser.get<std::vector<std::string>>("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<std::string> distributions;
{
auto ds = parser.get<std::vector<std::string>>("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<hiprandStateXORWOW_t>(parser, distribution);
run_benchmarks<curandStateXORWOW_t>(parser, distribution);
}
else if (engine == "mrg32k3a")
{
// CHECK: run_benchmarks<hiprandStateMRG32k3a_t>(parser, distribution);
run_benchmarks<curandStateMRG32k3a_t>(parser, distribution);
}
else if (engine == "philox")
{
// CHECK: run_benchmarks<hiprandStatePhilox4_32_10_t>(parser, distribution);
run_benchmarks<curandStatePhilox4_32_10_t>(parser, distribution);
}
else if (engine == "sobol32")
{
// CHECK: run_benchmarks<hiprandStateSobol32_t>(parser, distribution);
run_benchmarks<curandStateSobol32_t>(parser, distribution);
}
else if (engine == "mtgp32")
{
// CHECK: run_benchmarks<hiprandStateMtgp32_t>(parser, distribution);
run_benchmarks<curandStateMtgp32_t>(parser, distribution);
}
}
}
return 0;
}
projects/hip/tests/hipify-clang/cuRAND/cmdparser.hpp
+513
Просмотреть файл
@@ -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 <iostream>
#include <stdexcept>
#include <string>
#include <vector>
#include <sstream>
#include <functional>
namespace cli {
struct CallbackArgs {
const std::vector<std::string>& 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<std::string> 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<typename T>
struct ArgumentCountChecker
{
static constexpr bool Variadic = false;
};
template<typename T>
struct ArgumentCountChecker<std::vector<T>>
{
static constexpr bool Variadic = true;
};
template<typename T>
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<T>::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<T(CallbackArgs&)> callback;
T value;
};
template<typename T>
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<T>::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<std::string>& elements, const int&) {
if (elements.size() != 1)
throw std::bad_cast();
return std::stoi(elements[0]);
}
static bool parse(const std::vector<std::string>& 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<std::string>& elements, const double&) {
if (elements.size() != 1)
throw std::bad_cast();
return std::stod(elements[0]);
}
static float parse(const std::vector<std::string>& elements, const float&) {
if (elements.size() != 1)
throw std::bad_cast();
return std::stof(elements[0]);
}
static long double parse(const std::vector<std::string>& elements, const long double&) {
if (elements.size() != 1)
throw std::bad_cast();
return std::stold(elements[0]);
}
static unsigned int parse(const std::vector<std::string>& elements, const unsigned int&) {
if (elements.size() != 1)
throw std::bad_cast();
return static_cast<unsigned int>(std::stoul(elements[0]));
}
static unsigned long parse(const std::vector<std::string>& 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<std::string>& elements, const unsigned long long&) {
if (elements.size() != 1)
throw std::bad_cast();
return std::stoull(elements[0]);
}
static long parse(const std::vector<std::string>& elements, const long&) {
if (elements.size() != 1)
throw std::bad_cast();
return std::stol(elements[0]);
}
static std::string parse(const std::vector<std::string>& elements, const std::string&) {
if (elements.size() != 1)
throw std::bad_cast();
return elements[0];
}
template<class T>
static std::vector<T> parse(const std::vector<std::string>& elements, const std::vector<T>&) {
const T defval = T();
std::vector<T> values { };
std::vector<std::string> buffer(1);
for (const auto& element : elements) {
buffer[0] = element;
values.push_back(parse(buffer, defval));
}
return values;
}
template<class T>
static std::string stringify(const T& value) {
return std::to_string(value);
}
template<class T>
static std::string stringify(const std::vector<T>& 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<bool(CallbackArgs&)>([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<typename T>
void set_default(bool is_required, const std::string& description = "") {
auto command = new CmdArgument<T> { "", "", description, is_required, false };
_commands.push_back(command);
}
template<typename T>
void set_required(const std::string& name, const std::string& alternative, const std::string& description = "", bool dominant = false) {
auto command = new CmdArgument<T> { name, alternative, description, true, dominant };
_commands.push_back(command);
}
template<typename T>
void set_optional(const std::string& name, const std::string& alternative, T defaultValue, const std::string& description = "", bool dominant = false) {
auto command = new CmdArgument<T> { name, alternative, description, false, dominant };
command->value = defaultValue;
_commands.push_back(command);
}
template<typename T>
void set_callback(const std::string& name, const std::string& alternative, std::function<T(CallbackArgs&)> callback, const std::string& description = "", bool dominant = false) {
auto command = new CmdFunction<T> { 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<typename T>
T get(const std::string& name) const {
for (const auto& command : _commands) {
if (command->name == name) {
auto cmd = dynamic_cast<CmdArgument<T>*>(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<typename T>
T get_if(const std::string& name, std::function<T(T)> callback) const {
auto value = get<T>(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<int>(_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<std::string> _arguments;
std::vector<CmdBase*> _commands;
};
}
projects/hip/tests/hipify-clang/headers_test_09.cu
+93
Просмотреть файл
@@ -0,0 +1,93 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// CHECK: #include <hip/hip_runtime.h>
// CHECK: #include <memory>
// CHECK-NOT: #include <cuda_runtime.h>
// CHECK-NOT: #include <hip/hip_runtime.h>
// 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 <iostream>
// CHECK: #include "hipblas.h"
// CHECK-NOT: #include "cublas.h"
// CHECK: #include <stdio.h>
// CHECK: #include "hiprand.h"
// CHECK: #include "hiprand_kernel.h"
// CHECK: #include <algorithm>
// 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 <string>
#include <cuda.h>
#include <memory>
#include <cuda_runtime.h>
#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 <iostream>
#include "cublas_v2.h"
#include "cublas.h"
#include <stdio.h>
#include "curand.h"
#include "curand_kernel.h"
#include <algorithm>
#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 <string>
projects/hip/tests/hipify-clang/lit.cfg
+2
Просмотреть файл
@@ -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