From 290b5f17382b6fd192aed06de6095621c4c97fef Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 19 Dec 2017 16:06:14 +0530 Subject: [PATCH 01/10] Implement hipStreamAddCallback Change-Id: Ib851e4d86ba9c8406ca37b88162ea483ccbc9d36 [ROCm/clr commit: cd9ba0d1e18ca6bc8242357e8d10346b324e27a3] --- projects/clr/hipamd/src/hip_hcc.cpp | 30 +++++++++ projects/clr/hipamd/src/hip_hcc_internal.h | 20 ++++++ projects/clr/hipamd/src/hip_stream.cpp | 36 +++++----- .../stream/hipStreamAddCallback.cpp | 66 +++++++++++++------ 4 files changed, 114 insertions(+), 38 deletions(-) diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index f15a0eb1d8..0ca170152b 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -37,6 +37,7 @@ THE SOFTWARE. #include #include #include +#include #include #include @@ -1409,9 +1410,38 @@ void ihipInit() tprintf(DB_SYNC, "pid=%u %-30s g_numLogicalThreads=%u\n", getpid(), "", g_numLogicalThreads); } +hipError_t ihipStreamSynchronize(hipStream_t stream) +{ + hipError_t e = hipSuccess; + if (stream == hipStreamNull) { + ihipCtx_t *ctx = ihipGetTlsDefaultCtx(); + ctx->locked_syncDefaultStream(true/*waitOnSelf*/, true/*syncToHost*/); + } else { + // note this does not synchornize with the NULL stream: + stream->locked_wait(); + e = hipSuccess; + } + return e; +} +void ihipStreamCallbackHandler(ihipStreamCallback_t *cb) +{ + hipError_t e = hipSuccess; + + // Notify hipStreamAddCallback that callback handler thread is active + std::lock_guard guard(cb->_mtx); + cb->_ready = true; + + // Synchronize stream + tprintf(DB_SYNC, "ihipStreamCallbackHandler wait on stream %s\n", ToString(cb->_stream).c_str()); + e = ihipStreamSynchronize(cb->_stream); + + // Call registered callback function + cb->_callback(cb->_stream, e, cb->_userData); + delete cb; +} //--- // Get the stream to use for a command submission. diff --git a/projects/clr/hipamd/src/hip_hcc_internal.h b/projects/clr/hipamd/src/hip_hcc_internal.h index 4891f54fee..601b66f343 100644 --- a/projects/clr/hipamd/src/hip_hcc_internal.h +++ b/projects/clr/hipamd/src/hip_hcc_internal.h @@ -622,6 +622,24 @@ private: // Data }; +//---- +// Internal structure for stream callback handler +class ihipStreamCallback_t { +public: + ihipStreamCallback_t(hipStream_t stream, hipStreamCallback_t callback, void *userData) : + _stream(stream), + _callback(callback), + _userData(userData) + { + _ready = false; + }; + hipStream_t _stream; + hipStreamCallback_t _callback; + void* _userData; + std::mutex _mtx; + bool _ready; +}; + //---- // Internal event structure: @@ -931,6 +949,8 @@ ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex); hipStream_t ihipSyncAndResolveStream(hipStream_t); +hipError_t ihipStreamSynchronize(hipStream_t stream); +void ihipStreamCallbackHandler(ihipStreamCallback_t *cb); // Stream printf functions: inline std::ostream& operator<<(std::ostream& os, const ihipStream_t& s) diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index dab31dad62..94fc436b75 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -20,6 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +#include +#include #include "hip/hip_runtime.h" #include "hip_hcc_internal.h" #include "trace_helper.h" @@ -147,20 +149,8 @@ hipError_t hipStreamSynchronize(hipStream_t stream) { HIP_INIT_SPECIAL_API(TRACE_SYNC, stream); - hipError_t e = hipSuccess; - - if (stream == hipStreamNull) { - ihipCtx_t *ctx = ihipGetTlsDefaultCtx(); - ctx->locked_syncDefaultStream(true/*waitOnSelf*/, true/*syncToHost*/); - } else { - // note this does not synchornize with the NULL stream: - stream->locked_wait(); - e = hipSuccess; - } - - - return ihipLogStatus(e); -}; + return ihipLogStatus(ihipStreamSynchronize(stream)); +} //--- @@ -216,8 +206,20 @@ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback { HIP_INIT_API(stream, callback, userData, flags); hipError_t e = hipSuccess; - //--- explicitly synchronize stream to add callback routines - hipStreamSynchronize(stream); - callback(stream, e, userData); + + // Create a thread in detached mode to handle callback + ihipStreamCallback_t *cb = new ihipStreamCallback_t(stream, callback, userData); + std::thread (ihipStreamCallbackHandler, cb).detach(); + + // Wait for thread to be ready + cb->_mtx.lock(); + while(cb->_ready != true) + { + cb->_mtx.unlock(); + std::this_thread::sleep_for(std::chrono::milliseconds(10)); + cb->_mtx.lock(); + } + cb->_mtx.unlock(); + return ihipLogStatus(e); } diff --git a/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp b/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp index 32a2793479..692d090509 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp @@ -23,8 +23,7 @@ THE SOFTWARE. * HIT_END */ -// Test under-development. Call hipStreamAddCallback function and see if it works as expected. - +#include #include "hip/hip_runtime.h" #include "test_common.h" @@ -32,32 +31,57 @@ THE SOFTWARE. #define HIPRT_CB #endif -class CallbackClass +__global__ void vector_square(float *C_d, float *A_d, size_t N) { -public: - static void HIPRT_CB Callback(hipStream_t stream, hipError_t status, void *userData); + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x ; -private: - void callbackFunc(hipError_t status); -}; - -void HIPRT_CB CallbackClass::Callback(hipStream_t stream, hipError_t status, void *userData) -{ - CallbackClass* obj = (CallbackClass*) userData; - obj->callbackFunc(status); + for (size_t i=offset; i Date: Mon, 29 Jan 2018 18:33:47 +0300 Subject: [PATCH 02/10] [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/clr commit: 58f2d3c93fcd7012025a761060e9effe1531e866] --- .../markdown/CURAND_API_supported_by_HIP.md | 2 + .../hipamd/hipify-clang/src/CUDA2HipMap.cpp | 45 +- .../hipamd/hipify-clang/src/HipifyAction.cpp | 68 +- .../hipamd/hipify-clang/src/HipifyAction.h | 5 + .../hipamd/hipify-clang/src/Statistics.cpp | 4 +- .../clr/hipamd/hipify-clang/src/Statistics.h | 4 +- .../cuRAND/benchmark_curand_generate.cpp | 393 ++++++++++ .../cuRAND/benchmark_curand_kernel.cpp | 669 ++++++++++++++++++ .../tests/hipify-clang/cuRAND/cmdparser.hpp | 513 ++++++++++++++ .../tests/hipify-clang/headers_test_09.cu | 93 +++ .../clr/hipamd/tests/hipify-clang/lit.cfg | 2 + 11 files changed, 1756 insertions(+), 42 deletions(-) create mode 100644 projects/clr/hipamd/tests/hipify-clang/cuRAND/benchmark_curand_generate.cpp create mode 100644 projects/clr/hipamd/tests/hipify-clang/cuRAND/benchmark_curand_kernel.cpp create mode 100644 projects/clr/hipamd/tests/hipify-clang/cuRAND/cmdparser.hpp create mode 100644 projects/clr/hipamd/tests/hipify-clang/headers_test_09.cu diff --git a/projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md b/projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md index 900cfadc03..aae43af272 100644 --- a/projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md +++ b/projects/clr/hipamd/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/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp index 6bcc0b38c4..7001deb2e7 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/clr/hipamd/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/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp index 3459ea21f9..76efe9bf85 100644 --- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/projects/clr/hipamd/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/projects/clr/hipamd/hipify-clang/src/HipifyAction.h b/projects/clr/hipamd/hipify-clang/src/HipifyAction.h index 42622c1e01..8c2195b28d 100644 --- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.h +++ b/projects/clr/hipamd/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/projects/clr/hipamd/hipify-clang/src/Statistics.cpp b/projects/clr/hipamd/hipify-clang/src/Statistics.cpp index e2b3632066..4c5664a6b0 100644 --- a/projects/clr/hipamd/hipify-clang/src/Statistics.cpp +++ b/projects/clr/hipamd/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/projects/clr/hipamd/hipify-clang/src/Statistics.h b/projects/clr/hipamd/hipify-clang/src/Statistics.h index f160ca7383..81be7b09a8 100644 --- a/projects/clr/hipamd/hipify-clang/src/Statistics.h +++ b/projects/clr/hipamd/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/projects/clr/hipamd/tests/hipify-clang/cuRAND/benchmark_curand_generate.cpp b/projects/clr/hipamd/tests/hipify-clang/cuRAND/benchmark_curand_generate.cpp new file mode 100644 index 0000000000..8d284c82c6 --- /dev/null +++ b/projects/clr/hipamd/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/projects/clr/hipamd/tests/hipify-clang/cuRAND/benchmark_curand_kernel.cpp b/projects/clr/hipamd/tests/hipify-clang/cuRAND/benchmark_curand_kernel.cpp new file mode 100644 index 0000000000..2c77f25331 --- /dev/null +++ b/projects/clr/hipamd/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/projects/clr/hipamd/tests/hipify-clang/cuRAND/cmdparser.hpp b/projects/clr/hipamd/tests/hipify-clang/cuRAND/cmdparser.hpp new file mode 100644 index 0000000000..364612ebb1 --- /dev/null +++ b/projects/clr/hipamd/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/projects/clr/hipamd/tests/hipify-clang/headers_test_09.cu b/projects/clr/hipamd/tests/hipify-clang/headers_test_09.cu new file mode 100644 index 0000000000..048ac2e9a9 --- /dev/null +++ b/projects/clr/hipamd/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/projects/clr/hipamd/tests/hipify-clang/lit.cfg b/projects/clr/hipamd/tests/hipify-clang/lit.cfg index 12b1410cee..104cee8311 100644 --- a/projects/clr/hipamd/tests/hipify-clang/lit.cfg +++ b/projects/clr/hipamd/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 From b2bb274b7e06290e01e841807d54a2cd7ebd4dfa Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 30 Jan 2018 12:32:46 +0530 Subject: [PATCH 03/10] Fix hipStreamAddCallback testcase for nvcc Change-Id: Ieec4b8d7933d8d68394d21d27132da206111efc8 [ROCm/clr commit: 6a37c60eaa11832d2c99b63465645fa090e6ce3c] --- .../tests/src/runtimeApi/stream/hipStreamAddCallback.cpp | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp b/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp index 692d090509..0f9f32a253 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp @@ -24,6 +24,7 @@ THE SOFTWARE. */ #include +#include #include "hip/hip_runtime.h" #include "test_common.h" @@ -42,16 +43,17 @@ __global__ void vector_square(float *C_d, float *A_d, size_t N) } float *A_h, *C_h; -size_t N = 1000000; +bool cbDone = false; static void HIPRT_CB Callback(hipStream_t stream, hipError_t status, void *userData) { for (size_t i=0; i Date: Tue, 30 Jan 2018 18:55:09 +0300 Subject: [PATCH 04/10] [HIPIFY][tests] fix typo [ROCm/clr commit: 2ae2693c8a359d03d9a199d436155ee8e9e04d56] --- projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp | 2 +- .../tests/hipify-clang/cuRAND/benchmark_curand_kernel.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp index 7001deb2e7..35dab03430 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp @@ -2872,7 +2872,7 @@ const std::map CUDA_IDENTIFIER_MAP{ {"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}}, - {"curand_discrete", {"hiprand_discrete4", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_discrete", {"hiprand_discrete", CONV_DEVICE_FUNC, API_RAND}}, {"curand_discrete4", {"hiprand_discrete4", CONV_DEVICE_FUNC, API_RAND}}, {"curand_poisson", {"hiprand_poisson", CONV_DEVICE_FUNC, API_RAND}}, {"curand_poisson4", {"hiprand_poisson4", CONV_DEVICE_FUNC, API_RAND}}, diff --git a/projects/clr/hipamd/tests/hipify-clang/cuRAND/benchmark_curand_kernel.cpp b/projects/clr/hipamd/tests/hipify-clang/cuRAND/benchmark_curand_kernel.cpp index 2c77f25331..bdcf512993 100644 --- a/projects/clr/hipamd/tests/hipify-clang/cuRAND/benchmark_curand_kernel.cpp +++ b/projects/clr/hipamd/tests/hipify-clang/cuRAND/benchmark_curand_kernel.cpp @@ -510,7 +510,7 @@ void run_benchmarks(const cli::Parser& parser, 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); + // CHECK: return hiprand_discrete(state, discrete_distribution); return curand_discrete(state, discrete_distribution); }, discrete_distribution ); From 2a1267bd43b85f85d70b6793a9019064a84558ba Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 30 Jan 2018 19:50:18 +0300 Subject: [PATCH 05/10] [HIPIFY][tests] add poisson-api-example test + add missing types + doc update [ROCm/clr commit: eef94a74bc48a27941570fd08837ad2b386bde0b] --- .../markdown/CURAND_API_supported_by_HIP.md | 2 + .../hipamd/hipify-clang/src/CUDA2HipMap.cpp | 4 +- .../cuRAND/poisson_api_example.cu | 417 ++++++++++++++++++ 3 files changed, 422 insertions(+), 1 deletion(-) create mode 100644 projects/clr/hipamd/tests/hipify-clang/cuRAND/poisson_api_example.cu diff --git a/projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md b/projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md index aae43af272..c35eeb26b6 100644 --- a/projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md +++ b/projects/clr/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md @@ -101,6 +101,8 @@ | struct | `curandStateMRG32k3a_t` | `hiprandStateMRG32k3a_t` | | struct | `curandStatePhilox4_32_10_t` | `hiprandStatePhilox4_32_10_t` | | struct | `curandStateXORWOW_t` | `hiprandStateXORWOW_t` | +| struct | `curandState_t` | `hiprandState_t` | +| struct | `curandState` | `hiprandState_t` | ## **2. Host API Functions** diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp index 35dab03430..47358802e9 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HipMap.cpp @@ -362,6 +362,8 @@ const std::map CUDA_TYPE_NAME_MAP{ {"curandStateMRG32k3a_t", {"hiprandStateMRG32k3a_t", CONV_TYPE, API_RAND}}, {"curandStatePhilox4_32_10_t", {"hiprandStatePhilox4_32_10_t", CONV_TYPE, API_RAND}}, {"curandStateXORWOW_t", {"hiprandStateXORWOW_t", CONV_TYPE, API_RAND}}, + {"curandState_t", {"hiprandState_t", CONV_TYPE, API_RAND}}, + {"curandState", {"hiprandState_t", CONV_TYPE, API_RAND}}, }; /// Maps cuda header names to hip header names. @@ -2872,7 +2874,7 @@ const std::map CUDA_IDENTIFIER_MAP{ {"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}}, - {"curand_discrete", {"hiprand_discrete", CONV_DEVICE_FUNC, API_RAND}}, + {"curand_discrete", {"hiprand_discrete", CONV_DEVICE_FUNC, API_RAND}}, {"curand_discrete4", {"hiprand_discrete4", CONV_DEVICE_FUNC, API_RAND}}, {"curand_poisson", {"hiprand_poisson", CONV_DEVICE_FUNC, API_RAND}}, {"curand_poisson4", {"hiprand_poisson4", CONV_DEVICE_FUNC, API_RAND}}, diff --git a/projects/clr/hipamd/tests/hipify-clang/cuRAND/poisson_api_example.cu b/projects/clr/hipamd/tests/hipify-clang/cuRAND/poisson_api_example.cu new file mode 100644 index 0000000000..d4cfd90e1f --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/cuRAND/poisson_api_example.cu @@ -0,0 +1,417 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// Taken from: http://docs.nvidia.com/cuda/curand/device-api-overview.html#poisson-api-example +/* + * This program uses CURAND library for Poisson distribution + * to simulate queues in store for 16 hours. It shows the + * difference of using 3 different APIs: + * - HOST API -arrival of customers is described by Poisson(4) + * - SIMPLE DEVICE API -arrival of customers is described by + * Poisson(4*(sin(x/100)+1)), where x is number of minutes + * from store opening time. + * - ROBUST DEVICE API -arrival of customers is described by: + * - Poisson(2) for first 3 hours. + * - Poisson(1) for second 3 hours. + * - Poisson(3) after 6 hours. + */ + +#include +#include +// CHECK: #include +#include +// CHECK: #include +#include +// CHECK: #include +#include + +// CHECK: #define CUDA_CALL(x) do { if((x) != hipSuccess) { +#define CUDA_CALL(x) do { if((x) != cudaSuccess) { \ + printf("Error at %s:%d\n",__FILE__,__LINE__); \ + return EXIT_FAILURE;}} while(0) +// CHECK: #define CURAND_CALL(x) do { if((x)!=HIPRAND_STATUS_SUCCESS) { +#define CURAND_CALL(x) do { if((x)!=CURAND_STATUS_SUCCESS) { \ + printf("Error at %s:%d\n",__FILE__,__LINE__);\ + return EXIT_FAILURE;}} while(0) + + +#define HOURS 16 +#define OPENING_HOUR 7 +#define CLOSING_HOUR (OPENING_HOUR + HOURS) + +#define access_2D(type, ptr, row, column, pitch)\ + *((type*)((char*)ptr + (row) * pitch) + column) + +enum API_TYPE { + HOST_API = 0, + SIMPLE_DEVICE_API = 1, + ROBUST_DEVICE_API = 2, +}; + +/* global variables */ +API_TYPE api; +int report_break; +int cashiers_load_h[HOURS]; +__constant__ int cashiers_load[HOURS]; +// CHECK: __global__ void setup_kernel(hiprandState_t *state) +__global__ void setup_kernel(curandState *state) +{ + int id = threadIdx.x + blockIdx.x * blockDim.x; + /* Each thread gets same seed, a different sequence + number, no offset */ + // CHECK: hiprand_init(1234, id, 0, &state[id]); + curand_init(1234, id, 0, &state[id]); +} + +__inline__ __device__ +void update_queue(int id, int min, unsigned int new_customers, + unsigned int &queue_length, + unsigned int *queue_lengths, size_t pitch) +{ + int balance; + balance = new_customers - 2 * cashiers_load[(min-1)/60]; + if (balance + (int)queue_length <= 0){ + queue_length = 0; + }else{ + queue_length += balance; + } + /* Store results */ + access_2D(unsigned int, queue_lengths, min-1, id, pitch) + = queue_length; +} + +// CHECK: __global__ void simple_device_API_kernel(hiprandState_t *state, +__global__ void simple_device_API_kernel(curandState *state, + unsigned int *queue_lengths, size_t pitch) +{ + int id = threadIdx.x + blockIdx.x * blockDim.x; + unsigned int new_customers; + unsigned int queue_length = 0; + /* Copy state to local memory for efficiency */ + // CHECK: hiprandState_t localState = state[id]; + curandState localState = state[id]; + /* Simulate queue in time */ + for(int min = 1; min <= 60 * HOURS; min++) { + /* Draw number of new customers depending on API */ + // CHECK: new_customers = hiprand_poisson(&localState, + new_customers = curand_poisson(&localState, + 4*(sin((float)min/100.0)+1)); + /* Update queue */ + update_queue(id, min, new_customers, queue_length, + queue_lengths, pitch); + } + /* Copy state back to global memory */ + state[id] = localState; +} + + +__global__ void host_API_kernel(unsigned int *poisson_numbers, + unsigned int *queue_lengths, size_t pitch) +{ + int id = threadIdx.x + blockIdx.x * blockDim.x; + unsigned int new_customers; + unsigned int queue_length = 0; + /* Simulate queue in time */ + for(int min = 1; min <= 60 * HOURS; min++) { + /* Get random number from global memory */ + new_customers = poisson_numbers + [blockDim.x * gridDim.x * (min -1) + id]; + /* Update queue */ + update_queue(id, min, new_customers, queue_length, + queue_lengths, pitch); + } +} +// CHECK: __global__ void robust_device_API_kernel(hiprandState_t *state, +// CHECK: hiprandDiscreteDistribution_t poisson_1, +// CHECK: hiprandDiscreteDistribution_t poisson_2, +// CHECK: hiprandDiscreteDistribution_t poisson_3, +__global__ void robust_device_API_kernel(curandState *state, + curandDiscreteDistribution_t poisson_1, + curandDiscreteDistribution_t poisson_2, + curandDiscreteDistribution_t poisson_3, + unsigned int *queue_lengths, size_t pitch) +{ + int id = threadIdx.x + blockIdx.x * 64; + unsigned int new_customers; + unsigned int queue_length = 0; + /* Copy state to local memory for efficiency */ + // CHECK: hiprandState_t localState = state[id]; + curandState localState = state[id]; + /* Simulate queue in time */ + /* first 3 hours */ + for(int min = 1; min <= 60 * 3; min++) { + /* draw number of new customers depending on API */ + new_customers = + // CHECK: hiprand_discrete(&localState, poisson_2); + curand_discrete(&localState, poisson_2); + /* Update queue */ + update_queue(id, min, new_customers, queue_length, + queue_lengths, pitch); + } + /* second 3 hours */ + for(int min = 60 * 3 + 1; min <= 60 * 6; min++) { + /* draw number of new customers depending on API */ + new_customers = + // CHECK: hiprand_discrete(&localState, poisson_1); + curand_discrete(&localState, poisson_1); + /* Update queue */ + update_queue(id, min, new_customers, queue_length, + queue_lengths, pitch); + } + /* after 6 hours */ + for(int min = 60 * 6 + 1; min <= 60 * HOURS; min++) { + /* draw number of new customers depending on API */ + new_customers = + // CHECK: hiprand_discrete(&localState, poisson_3); + curand_discrete(&localState, poisson_3); + /* Update queue */ + update_queue(id, min, new_customers, queue_length, + queue_lengths, pitch); + } + /* Copy state back to global memory */ + state[id] = localState; +} + +/* Set time intervals between reports */ +void report_settings() +{ + do{ + printf("Set time intervals between queue reports"); + printf("(in minutes > 0)\n"); + if (scanf("%d", &report_break) == 0) continue; + }while(report_break <= 0); +} + + +/* Set number of cashiers each hour */ +void add_cachiers(int *cashiers_load) +{ + int i, min, max, begin, end; + printf("Cashier serves 2 customers per minute...\n"); + for (i = 0; i < HOURS; i++){ + cashiers_load_h[i] = 0; + } + while (true){ + printf("Adding cashier...\n"); + min = OPENING_HOUR; + max = CLOSING_HOUR-1; + do{ + printf("Set hour that cahier comes (%d-%d)", + min, max); + printf(" [type 0 to finish adding cashiers]\n"); + if (scanf("%d", &begin) == 0) continue; + }while (begin > max || (begin < min && begin != 0)); + if (begin == 0) break; + min = begin+1; + max = CLOSING_HOUR; + do{ + printf("Set hour that cahier leaves (%d-%d)", + min, max); + printf(" [type 0 to finish adding cashiers]\n"); + if (scanf("%d", &end) == 0) continue; + }while (end > max || (end < min && end != 0)); + if (end == 0) break; + for (i = begin - OPENING_HOUR; + i < end - OPENING_HOUR; i++){ + cashiers_load_h[i]++; + } + } + for (i = OPENING_HOUR; i < CLOSING_HOUR; i++){ + printf("\n%2d:00 - %2d:00 %d cashier", + i, i+1, cashiers_load_h[i-OPENING_HOUR]); + if (cashiers_load[i-OPENING_HOUR] != 1) printf("s"); + } + printf("\n"); +} + +/* Set API type */ +API_TYPE set_API_type() +{ + printf("Choose API type:\n"); + int choose; + do{ + printf("type 1 for HOST API\n"); + printf("type 2 for SIMPLE DEVICE API\n"); + printf("type 3 for ROBUST DEVICE API\n"); + if (scanf("%d", &choose) == 0) continue; + }while( choose < 1 || choose > 3); + switch(choose){ + case 1: return HOST_API; + case 2: return SIMPLE_DEVICE_API; + case 3: return ROBUST_DEVICE_API; + default: + fprintf(stderr, "wrong API\n"); + return HOST_API; + } +} + +void settings() +{ + add_cachiers(cashiers_load); + // CHECK: hipMemcpyToSymbol("cashiers_load", cashiers_load_h, + // CHECK: HOURS * sizeof(int), 0, hipMemcpyHostToDevice); + cudaMemcpyToSymbol("cashiers_load", cashiers_load_h, + HOURS * sizeof(int), 0, cudaMemcpyHostToDevice); + report_settings(); + api = set_API_type(); +} + +void print_statistics(unsigned int *hostResults, size_t pitch) +{ + int min, i, hour, minute; + unsigned int sum; + for(min = report_break; min <= 60 * HOURS; + min += report_break) { + sum = 0; + for(i = 0; i < 64 * 64; i++) { + sum += access_2D(unsigned int, hostResults, + min-1, i, pitch); + } + hour = OPENING_HOUR + min/60; + minute = min%60; + printf("%2d:%02d # of waiting customers = %10.4g |", + hour, minute, (float)sum/(64.0 * 64.0)); + printf(" # of cashiers = %d | ", + cashiers_load_h[(min-1)/60]); + printf("# of new customers/min ~= "); + switch (api){ + case HOST_API: + printf("%2.2f\n", 4.0); + break; + case SIMPLE_DEVICE_API: + printf("%2.2f\n", + 4*(sin((float)min/100.0)+1)); + break; + case ROBUST_DEVICE_API: + if (min <= 3 * 60){ + printf("%2.2f\n", 2.0); + }else{ + if (min <= 6 * 60){ + printf("%2.2f\n", 1.0); + }else{ + printf("%2.2f\n", 3.0); + } + } + break; + default: + fprintf(stderr, "Wrong API\n"); + } + } +} + + +int main(int argc, char *argv[]) +{ + int n; + size_t pitch; + // CHECK: hiprandState_t *devStates; + curandState *devStates; + unsigned int *devResults, *hostResults; + unsigned int *poisson_numbers_d; + // CHECK: hiprandDiscreteDistribution_t poisson_1, poisson_2; + // CHECK: hiprandDiscreteDistribution_t poisson_3; + // CHECK: hiprandGenerator_t gen; + curandDiscreteDistribution_t poisson_1, poisson_2; + curandDiscreteDistribution_t poisson_3; + curandGenerator_t gen; + + /* Setting cashiers, report and API */ + settings(); + + /* Allocate space for results on device */ + // CHECK: CUDA_CALL(hipMallocPitch((void **)&devResults, &pitch, + CUDA_CALL(cudaMallocPitch((void **)&devResults, &pitch, + 64 * 64 * sizeof(unsigned int), 60 * HOURS)); + + /* Allocate space for results on host */ + hostResults = (unsigned int *)calloc(pitch * 60 * HOURS, + sizeof(unsigned int)); + + /* Allocate space for prng states on device */ + // CHECK: CUDA_CALL(hipMalloc((void **)&devStates, 64 * 64 * + // CHECK: sizeof(hiprandState_t))); + CUDA_CALL(cudaMalloc((void **)&devStates, 64 * 64 * + sizeof(curandState))); + + /* Setup prng states */ + if (api != HOST_API){ + // CHECK: hipLaunchKernelGGL(setup_kernel, dim3(64), dim3(64), 0, 0, devStates); + setup_kernel<<<64, 64>>>(devStates); + } + /* Simulate queue */ + switch (api){ + case HOST_API: + /* Create pseudo-random number generator */ + // CHECK: CURAND_CALL(hiprandCreateGenerator(&gen, + // CHECK: HIPRAND_RNG_PSEUDO_DEFAULT)); + CURAND_CALL(curandCreateGenerator(&gen, + CURAND_RNG_PSEUDO_DEFAULT)); + /* Set seed */ + // CHECK: CURAND_CALL(hiprandSetPseudoRandomGeneratorSeed( + CURAND_CALL(curandSetPseudoRandomGeneratorSeed( + gen, 1234ULL)); + /* compute n */ + n = 64 * 64 * HOURS * 60; + /* Allocate n unsigned ints on device */ + // CHECK: CUDA_CALL(hipMalloc((void **)&poisson_numbers_d, + CUDA_CALL(cudaMalloc((void **)&poisson_numbers_d, + n * sizeof(unsigned int))); + /* Generate n unsigned ints on device */ + // CHECK: CURAND_CALL(hiprandGeneratePoisson(gen, + CURAND_CALL(curandGeneratePoisson(gen, + poisson_numbers_d, n, 4.0)); + // CHECK: hipLaunchKernelGGL(host_API_kernel, dim3(64), dim3(64), 0, 0, poisson_numbers_d, + host_API_kernel<<<64, 64>>>(poisson_numbers_d, + devResults, pitch); + /* Cleanup */ + // CHECK: CURAND_CALL(hiprandDestroyGenerator(gen)); + CURAND_CALL(curandDestroyGenerator(gen)); + break; + case SIMPLE_DEVICE_API: + // CHECK: hipLaunchKernelGGL(simple_device_API_kernel, dim3(64), dim3(64), 0, 0, devStates, + simple_device_API_kernel<<<64, 64>>>(devStates, + devResults, pitch); + break; + case ROBUST_DEVICE_API: + /* Create histograms for Poisson(1) */ + // CHECK: CURAND_CALL(hiprandCreatePoissonDistribution(1.0, + CURAND_CALL(curandCreatePoissonDistribution(1.0, + &poisson_1)); + /* Create histograms for Poisson(2) */ + // CHECK: CURAND_CALL(hiprandCreatePoissonDistribution(2.0, + CURAND_CALL(curandCreatePoissonDistribution(2.0, + &poisson_2)); + /* Create histograms for Poisson(3) */ + // CHECK: CURAND_CALL(hiprandCreatePoissonDistribution(3.0, + CURAND_CALL(curandCreatePoissonDistribution(3.0, + &poisson_3)); + // CHECK: hipLaunchKernelGGL(robust_device_API_kernel, dim3(64), dim3(64), 0, 0, devStates, + robust_device_API_kernel<<<64, 64>>>(devStates, + poisson_1, poisson_2, poisson_3, + devResults, pitch); + /* Cleanup */ + // CHECK: CURAND_CALL(hiprandDestroyDistribution(poisson_1)); + // CHECK: CURAND_CALL(hiprandDestroyDistribution(poisson_2)); + // CHECK: CURAND_CALL(hiprandDestroyDistribution(poisson_3)); + CURAND_CALL(curandDestroyDistribution(poisson_1)); + CURAND_CALL(curandDestroyDistribution(poisson_2)); + CURAND_CALL(curandDestroyDistribution(poisson_3)); + break; + default: + fprintf(stderr, "Wrong API\n"); + } + /* Copy device memory to host */ + // CHECK: CUDA_CALL(hipMemcpy2D(hostResults, pitch, devResults, + // CHECK: 60 * HOURS, hipMemcpyDeviceToHost)); + CUDA_CALL(cudaMemcpy2D(hostResults, pitch, devResults, + pitch, 64 * 64 * sizeof(unsigned int), + 60 * HOURS, cudaMemcpyDeviceToHost)); + /* Show result */ + print_statistics(hostResults, pitch); + /* Cleanup */ + // CHECK: CUDA_CALL(hipFree(devStates)); + // CHECK: CUDA_CALL(hipFree(devResults)); + CUDA_CALL(cudaFree(devStates)); + CUDA_CALL(cudaFree(devResults)); + free(hostResults); + return EXIT_SUCCESS; +} From 469ffecc71fba2a238117e99dbdf0af5e6f2edc9 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 31 Jan 2018 10:48:17 +0530 Subject: [PATCH 06/10] hipStreamAddCallback.cpp: Replace unistd sleep with sleep_for Change-Id: I7a5d40a1acd8be76a0f175bcfa731ad89fb88d81 [ROCm/clr commit: 3aab52d7ccc37146a82642b03d0395c387bc2aa8] --- .../tests/src/runtimeApi/stream/hipStreamAddCallback.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp b/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp index 0f9f32a253..02912e14c6 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/stream/hipStreamAddCallback.cpp @@ -18,13 +18,14 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../../test_common.cpp + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 * RUN: %t * HIT_END */ #include -#include +#include +#include #include "hip/hip_runtime.h" #include "test_common.h" @@ -87,5 +88,6 @@ int main(int argc, char *argv[]) HIPCHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream)); HIPCHECK(hipStreamAddCallback(mystream, Callback, NULL, 0)); - while(!cbDone) sleep(1); + while(!cbDone) + std::this_thread::sleep_for(std::chrono::milliseconds(10)); } From 3976fbfe159d836559586c2afbe4fa38ee641aed Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 1 Feb 2018 17:07:48 +0300 Subject: [PATCH 07/10] [HIPIFY][tests] Add intro.cu test [ROCm/clr commit: 7d71dfd1f738a01188ee39dced233e8d281f9dd1] --- .../clr/hipamd/tests/hipify-clang/intro.cu | 174 ++++++++++++++++++ 1 file changed, 174 insertions(+) create mode 100644 projects/clr/hipamd/tests/hipify-clang/intro.cu diff --git a/projects/clr/hipamd/tests/hipify-clang/intro.cu b/projects/clr/hipamd/tests/hipify-clang/intro.cu new file mode 100644 index 0000000000..4b9c5c0da7 --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/intro.cu @@ -0,0 +1,174 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +#include +#include +#include +// CHECK: #include +#include + +#define K_THREADS 64 +#define K_INDEX() ((gridDim.x * blockIdx.y + blockIdx.x) * blockDim.x + threadIdx.x) +#define RND() ((rand() & 0x7FFF) / float(0x8000)) +#define ERRORCHECK() cErrorCheck(__FILE__, __LINE__) + +// CHECK: hipEvent_t t##_start, t##_end; \ +// CHECK: hipEventCreate(&t##_start); \ +// CHECK: hipEventCreate(&t##_end); +#define TIMER_CREATE(t) \ + cudaEvent_t t##_start, t##_end; \ + cudaEventCreate(&t##_start); \ + cudaEventCreate(&t##_end); + +// CHECK: hipEventRecord(t##_start); \ +// CHECK: hipEventSynchronize(t##_start); +#define TIMER_START(t) \ + cudaEventRecord(t##_start); \ + cudaEventSynchronize(t##_start); \ + +// CHECK: hipEventRecord(t##_start); \ +// CHECK: hipEventSynchronize(t##_start); \ +// CHECK: hipEventRecord(t##_end); \ +// CHECK: hipEventSynchronize(t##_end); \ +// CHECK: hipEventElapsedTime(&t, t##_start, t##_end); +#define TIMER_END(t) \ + cudaEventRecord(t##_start); \ + cudaEventSynchronize(t##_start); \ + cudaEventRecord(t##_end); \ + cudaEventSynchronize(t##_end); \ + cudaEventElapsedTime(&t, t##_start, t##_end); + + +inline void cErrorCheck(const char *file, int line) { +// CHECK: hipDeviceSynchronize(); +// CHECK: hipError_t err = hipGetLastError(); +// CHECK: if (err != hipSuccess) { +// CHECK: printf("Error: %s\n", hipGetErrorString(err)); + cudaThreadSynchronize(); + cudaError_t err = cudaGetLastError(); + if (err != cudaSuccess) { + printf("Error: %s\n", cudaGetErrorString(err)); + printf(" @ %s: %d\n", file, line); + exit(-1); + } +} + +inline dim3 K_GRID(int n, int threads = K_THREADS) { + int blocks = (int)ceilf(sqrtf((float)n/threads)); + dim3 grid(blocks, blocks); + return grid; +} + +typedef struct data { + int n; + float4 *r, *v, *f; +} data; + +data cpu, gpu; + +#define N 20 + +__global__ void repulsion(data gpu); +__global__ void integration(data gpu); + + +int main() { + printf("Cuda Test 1\n"); + + int count = 0; + // CHECK: hipGetDeviceCount(&count); + cudaGetDeviceCount(&count); + printf(" %d CUDA devices found\n", count); + if(!count) { + ::exit(EXIT_FAILURE); + } + // CHECK: hipFree(0); + cudaFree(0); + + cpu.n = N; + + cpu.r = (float4*)malloc(N * sizeof(float4)); + cpu.v = (float4*)malloc(N * sizeof(float4)); + cpu.f = (float4*)malloc(N * sizeof(float4)); + + for(int i = 0; i < N; ++i) { + cpu.v[i] = make_float4(0,0,0,0); + cpu.r[i] = make_float4(RND(), RND(), RND(), 0); + cpu.f[i] = make_float4(0,0.01,0,0); + } + + gpu = cpu; + // CHECK: hipMalloc(&gpu.r, N * sizeof(float4)); + // CHECK: hipMalloc(&gpu.v, N * sizeof(float4)); + // CHECK: hipMalloc(&gpu.f, N * sizeof(float4)); + cudaMalloc(&gpu.r, N * sizeof(float4)); + cudaMalloc(&gpu.v, N * sizeof(float4)); + cudaMalloc(&gpu.f, N * sizeof(float4)); + // CHECK: hipMemcpy(gpu.r, cpu.r, cpu.n * sizeof(float4), hipMemcpyHostToDevice); + // CHECK: hipMemcpy(gpu.v, cpu.v, cpu.n * sizeof(float4), hipMemcpyHostToDevice); + // CHECK: hipMemcpy(gpu.f, cpu.f, cpu.n * sizeof(float4), hipMemcpyHostToDevice); + cudaMemcpy(gpu.r, cpu.r, cpu.n * sizeof(float4), cudaMemcpyHostToDevice); + cudaMemcpy(gpu.v, cpu.v, cpu.n * sizeof(float4), cudaMemcpyHostToDevice); + cudaMemcpy(gpu.f, cpu.f, cpu.n * sizeof(float4), cudaMemcpyHostToDevice); + + ERRORCHECK(); + float rep; + TIMER_CREATE(rep); + TIMER_START(rep); + // CHECK: hipLaunchKernelGGL(integration, dim3(K_GRID(cpu.n)), dim3(K_THREADS), 0, 0, gpu); + integration <<< K_GRID(cpu.n), K_THREADS >>>(gpu); + + TIMER_END(rep); + printf("Took: %f ms\n", rep); + ERRORCHECK(); + // CHECK: hipMemcpy(cpu.r, gpu.r, cpu.n * sizeof(float4), hipMemcpyDeviceToHost); + // CHECK: hipMemcpy(cpu.v, gpu.v, cpu.n * sizeof(float4), hipMemcpyDeviceToHost); + // CHECK: hipMemcpy(cpu.f, gpu.f, cpu.n * sizeof(float4), hipMemcpyDeviceToHost); + cudaMemcpy(cpu.r, gpu.r, cpu.n * sizeof(float4), cudaMemcpyDeviceToHost); + cudaMemcpy(cpu.v, gpu.v, cpu.n * sizeof(float4), cudaMemcpyDeviceToHost); + cudaMemcpy(cpu.f, gpu.f, cpu.n * sizeof(float4), cudaMemcpyDeviceToHost); + // CHECK: hipHostFree(cpu.r); + // CHECK: hipHostFree(cpu.v); + // CHECK: hipHostFree(cpu.f); + cudaFreeHost(cpu.r); + cudaFreeHost(cpu.v); + cudaFreeHost(cpu.f); + // CHECK: hipFree(gpu.r); + // CHECK: hipFree(gpu.v); + // CHECK: hipFree(gpu.f); + cudaFree(gpu.r); + cudaFree(gpu.v); + cudaFree(gpu.f); + // CHECK: hipDeviceReset(); + cudaDeviceReset(); + + printf("Results: \n"); + for(int i = 0; i < N; ++i) { + printf("%f, %f, %f \n", cpu.r[i].x, cpu.r[i].y, cpu.r[i].z); + } + + printf("Ready...\n"); + return 0; +} + +__global__ void repulsion(data gpu) { + int idx = K_INDEX(); + if(idx < N) { + gpu.r[idx].x = 1; + gpu.r[idx].y = 1; + gpu.r[idx].z = 1; + } +} + +#define MULT4(v, s) v.x *= s; v.y *= s; v.z *= s; v.w *= s; +#define ADD4(v1, v2) v1.x += v2.x; v1.y += v2.y; v1.z += v2.z; v1.w += v2.w; + +__global__ void integration(data gpu) { + int i = K_INDEX(); + if(i < N) { + MULT4(gpu.f[i], 0.01); + MULT4(gpu.v[i], 0.01); + ADD4(gpu.v[i], gpu.f[i]); + ADD4(gpu.r[i], gpu.v[i]); + gpu.f[i] = make_float4(0,0,0,0); + } +} \ No newline at end of file From d7d35ed98c47a552667ee6efee6d83496c4aeb16 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 1 Feb 2018 17:36:45 +0300 Subject: [PATCH 08/10] [HIPIFY][tests] add new line at the end of file [ROCm/clr commit: 91ef7b50e5d77dba7ffa52ad1f0a46cf467cb128] --- projects/clr/hipamd/tests/hipify-clang/intro.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/tests/hipify-clang/intro.cu b/projects/clr/hipamd/tests/hipify-clang/intro.cu index 4b9c5c0da7..da797eb2ec 100644 --- a/projects/clr/hipamd/tests/hipify-clang/intro.cu +++ b/projects/clr/hipamd/tests/hipify-clang/intro.cu @@ -171,4 +171,4 @@ __global__ void integration(data gpu) { ADD4(gpu.r[i], gpu.v[i]); gpu.f[i] = make_float4(0,0,0,0); } -} \ No newline at end of file +} From 180cf3f29ceb8eb8ac6a57fcbb1c12bcb4cb67ab Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 1 Feb 2018 18:34:16 +0300 Subject: [PATCH 09/10] [HIPIFY][tests] Add vec_add.cu test [ROCm/clr commit: f7fb6e2199b1b91f558b6162e869ae9bffa56c21] --- .../clr/hipamd/tests/hipify-clang/vec_add.cu | 90 +++++++++++++++++++ 1 file changed, 90 insertions(+) create mode 100644 projects/clr/hipamd/tests/hipify-clang/vec_add.cu diff --git a/projects/clr/hipamd/tests/hipify-clang/vec_add.cu b/projects/clr/hipamd/tests/hipify-clang/vec_add.cu new file mode 100644 index 0000000000..ec813e8bad --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/vec_add.cu @@ -0,0 +1,90 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// Kernel definition +__global__ void vecAdd(float* A, float* B, float* C) +{ + int i = threadIdx.x; + A[i] = 0; + B[i] = i; + C[i] = A[i] + B[i]; +} +// CHECK: #include +#include +#define SIZE 10 +#define KERNELINVOKES 5000000 +int vecadd(int gpudevice, int rank) +{ + int devcheck(int, int); + devcheck(gpudevice, rank); + float A[SIZE], B[SIZE], C[SIZE]; + // Kernel invocation + float *devPtrA; + float *devPtrB; + float *devPtrC; + int memsize = SIZE * sizeof(float); + // CHECK: hipMalloc((void**)&devPtrA, memsize); + // CHECK: hipMalloc((void**)&devPtrB, memsize); + // CHECK: hipMalloc((void**)&devPtrC, memsize); + cudaMalloc((void**)&devPtrA, memsize); + cudaMalloc((void**)&devPtrB, memsize); + cudaMalloc((void**)&devPtrC, memsize); + // CHECK: hipMemcpy(devPtrA, A, memsize, hipMemcpyHostToDevice); + // CHECK: hipMemcpy(devPtrB, B, memsize, hipMemcpyHostToDevice); + cudaMemcpy(devPtrA, A, memsize, cudaMemcpyHostToDevice); + cudaMemcpy(devPtrB, B, memsize, cudaMemcpyHostToDevice); + for (int i = 0; i>>(devPtrA, devPtrB, devPtrC); + } + // CHECK: hipMemcpy(C, devPtrC, memsize, hipMemcpyDeviceToHost); + cudaMemcpy(C, devPtrC, memsize, cudaMemcpyDeviceToHost); + // calculate only up to gpudevice to show the unique output + // of each rank's kernel launch + for (int i = 0; i= device_count) + { + printf("gpudevice >= device_count ... exiting\n"); + exit(1); + } + // CHECK: hipError_t cudareturn; + // CHECK: hipDeviceProp_t deviceProp; + // CHECK: hipGetDeviceProperties(&deviceProp, gpudevice); + cudaError_t cudareturn; + cudaDeviceProp deviceProp; + cudaGetDeviceProperties(&deviceProp, gpudevice); + // CHECK: if (deviceProp.hipWarpSize <= 1) + if (deviceProp.warpSize <= 1) + { + printf("rank %d: warning, CUDA Device Emulation (CPU) detected, exiting\n", rank); + exit(1); + } + // CHECK: cudareturn = hipSetDevice(gpudevice); + cudareturn = cudaSetDevice(gpudevice); + // CHECK: if (cudareturn == hipErrorInvalidDevice) + if (cudareturn == cudaErrorInvalidDevice) + { + // CHECK: perror("hipSetDevice returned hipErrorInvalidDevice"); + perror("cudaSetDevice returned cudaErrorInvalidDevice"); + } + else + { + // CHECK: hipGetDevice(&device); + cudaGetDevice(&device); + printf("rank %d: cudaGetDevice()=%d\n", rank, device); + } +} From a17cfac3ff6f753b830592aee9c34bd5ec0a6fe0 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 1 Feb 2018 19:41:36 +0300 Subject: [PATCH 10/10] [HIPIFY][tests] Add coalescing.cu test [ROCm/clr commit: e98ee1d62778d8b82d46b4c54f4cb9d1fc8bdcf0] --- .../hipamd/tests/hipify-clang/coalescing.cu | 117 ++++++++++++++++++ 1 file changed, 117 insertions(+) create mode 100644 projects/clr/hipamd/tests/hipify-clang/coalescing.cu diff --git a/projects/clr/hipamd/tests/hipify-clang/coalescing.cu b/projects/clr/hipamd/tests/hipify-clang/coalescing.cu new file mode 100644 index 0000000000..4c04289044 --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/coalescing.cu @@ -0,0 +1,117 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// To measure effects of memory coalescing. Coalescing.cu +// B. Wilkinson Jan 30, 2011 + +#include +#include +#include +// CHECK: #include +#include + +#define BlockSize 16 // Size of blocks, 32 x 32 threads, fixed, used globally + +__global__ void gpu_Comput (int *h, int N, int T) { + +// Array loaded with global thread ID that acesses that location + + int col = threadIdx.x + blockDim.x * blockIdx.x; + int row = threadIdx.y + blockDim.y * blockIdx.y; + + int threadID = col + row * N; + int index = row + col * N; // sequentially down each row + + for (int t = 0; t < T; t++) // loop to repeat to reduce other time effects + h[index] = threadID; // load array with flattened global thread ID +} + +void printArray(int *h, int N) { + + printf("Results of computation, every N/8 numbers, eight numbers\n"); + + for (int row = 0; row < N; row += N/8) { + for (int col = 0; col < N; col += N/8) + printf("%6d ", h[col + row * N]); + printf("\n"); + } +} + +int main(int argc, char *argv[]) { + + int T = 100; // number of iterations, entered at keyboard + int B = 1; // number of blocks, entered at keyboard + char key; + + int *h, *dev_h; // ptr to array holding numbers on host and device + // CHECK: hipEvent_t start, stop; + cudaEvent_t start, stop; // cuda events to measure time + float elapsed_time_ms1; + // CHECK: hipEventCreate( &start ); + // CHECK: hipEventCreate( &stop ); + cudaEventCreate( &start ); + cudaEventCreate( &stop ); + +/* ------------------------- Keyboard input -----------------------------------*/ + +do { // loop to repeat complete program + + printf("Grid Structure 2-D grid, 2-D blocks\n"); + printf("Blocks fixed at 16 x 16 threads, 512 threads, max for compute cap. 1.x\n"); + printf("Enter number of blocks in grid, each dimension, currently %d\n",B); + scanf("%d",&B); + printf("Enter number of iterations, currently %d\n",T); + scanf("%d",&T); + + int N = B * BlockSize; // size of data array, given input data + + printf("Array size (and total grid-block size) %d x %d\n", N, N); + + dim3 Block(BlockSize, BlockSize); //Block structure, 32 x 32 max + dim3 Grid(B, B); //Grid structure, B x B + +/* ------------------------- Allocate Memory-----------------------------------*/ + + int size = N * N * sizeof(int); // number of bytes in total in array + h = (int*) malloc(size); // Array on host + // CHECK: hipMalloc((void**)&dev_h, size); + cudaMalloc((void**)&dev_h, size); // allocate device memory + +/* ------------------------- GPU Computation -----------------------------------*/ + + // CHECK: hipEventRecord( start, 0 ); + cudaEventRecord( start, 0 ); + // CHECK: hipLaunchKernelGGL(gpu_Comput, dim3(Grid), dim3(Block), 0, 0, dev_h, N, T); + gpu_Comput<<< Grid, Block >>>(dev_h, N, T); + // CHECK: hipEventRecord( stop, 0 ); + // CHECK: hipEventSynchronize( stop ); + // CHECK: hipEventElapsedTime( &elapsed_time_ms1, start, stop ); + cudaEventRecord( stop, 0 ); // instrument code to measue end time + cudaEventSynchronize( stop ); // wait for all work done by threads + cudaEventElapsedTime( &elapsed_time_ms1, start, stop ); + // CHECK: hipMemcpy(h,dev_h, size ,hipMemcpyDeviceToHost); + cudaMemcpy(h,dev_h, size ,cudaMemcpyDeviceToHost); //Get results to check + + printArray(h,N); + printf("\nTime to calculate results on GPU: %f ms.\n", elapsed_time_ms1); + +/* -------------------------REPEAT PROGRAM INPUT-----------------------------------*/ + + printf("\nEnter c to repeat, return to terminate\n"); + + scanf("%c",&key); + scanf("%c",&key); + +} while (key == 'c'); // loop of complete program + +/* -------------- clean up ---------------------------------------*/ + +free(h); + // CHECK: hipFree(dev_h); + cudaFree(dev_h); + // CHECK: hipEventDestroy(start); + // CHECK: hipEventDestroy(stop); + cudaEventDestroy(start); + cudaEventDestroy(stop); + + return 0; +}