SWDEV-299127 - Merge 'develop' into 'amd-staging'
Change-Id: I6a2f7104f20e2dffcfa5b3f3c167f2cbf3488b0c
[ROCm/hip commit: dbb9b8476a]
Tento commit je obsažen v:
@@ -67,7 +67,7 @@ if(HIP_CXX_COMPILER MATCHES ".*hipcc")
|
||||
elseif (HIP_CXX_COMPILER MATCHES ".*clang\\+\\+")
|
||||
get_filename_component(_HIP_CLANG_REAL_PATH "${HIP_CXX_COMPILER}" REALPATH)
|
||||
get_filename_component(_HIP_CLANG_BIN_PATH "${_HIP_CLANG_REAL_PATH}" DIRECTORY)
|
||||
get_filename_component(HIP_CLANG_INSTALL_DIR "${_HIP_CLANG_BIN_PATH}" DIRECTORY)
|
||||
set(HIP_CLANG_PATH "${_HIP_CLANG_BIN_PATH}")
|
||||
endif()
|
||||
|
||||
|
||||
|
||||
+14
-17
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2023 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
|
||||
@@ -20,22 +20,19 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include "errorEnumerators.h"
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#ifndef HIP_INCLUDE_HIP_HIP_BF16_H
|
||||
#define HIP_INCLUDE_HIP_HIP_BF16_H
|
||||
|
||||
TEST_CASE("Unit_hipGetErrorString_Positive_Basic") {
|
||||
const char* error_string = nullptr;
|
||||
const auto enumerator =
|
||||
GENERATE(from_range(std::begin(kErrorEnumerators), std::end(kErrorEnumerators)));
|
||||
#include <hip/hip_common.h>
|
||||
|
||||
error_string = hipGetErrorString(enumerator);
|
||||
#if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && \
|
||||
!(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
|
||||
#include <hip/amd_detail/amd_hip_bf16.h>
|
||||
#elif !(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && \
|
||||
(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
|
||||
#include <hip/nvidia_detail/nvidia_hip_bf16.h>
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
|
||||
#endif
|
||||
|
||||
REQUIRE(error_string != nullptr);
|
||||
REQUIRE(strlen(error_string) > 0);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipGetErrorString_Negative_Parameters") {
|
||||
const char* error_string = hipGetErrorString(static_cast<hipError_t>(-1));
|
||||
REQUIRE(error_string != nullptr);
|
||||
}
|
||||
#endif // HIP_INCLUDE_HIP_HIP_BF16_H
|
||||
@@ -36,15 +36,15 @@ THE SOFTWARE.
|
||||
* to kernel params or extra arguments.
|
||||
*
|
||||
* @param [in] f Kernel to launch.
|
||||
* @param [in] gridDimX X grid dimension specified in work-items.
|
||||
* @param [in] gridDimY Y grid dimension specified in work-items.
|
||||
* @param [in] gridDimZ Z grid dimension specified in work-items.
|
||||
* @param [in] blockDimX X block dimension specified in work-items.
|
||||
* @param [in] blockDimY Y grid dimension specified in work-items.
|
||||
* @param [in] blockDimZ Z grid dimension specified in work-items.
|
||||
* @param [in] globalWorkSizeX X grid dimension specified in work-items.
|
||||
* @param [in] globalWorkSizeY Y grid dimension specified in work-items.
|
||||
* @param [in] globalWorkSizeZ Z grid dimension specified in work-items.
|
||||
* @param [in] localWorkSizeX X block dimension specified in work-items.
|
||||
* @param [in] localWorkSizeY Y block dimension specified in work-items.
|
||||
* @param [in] localWorkSizeZ Z block dimension specified in work-items.
|
||||
* @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel.
|
||||
* HIP-Clang compiler provides support for extern shared declarations.
|
||||
* @param [in] stream Stream where the kernel should be dispatched.
|
||||
* @param [in] hStream Stream where the kernel should be dispatched.
|
||||
* May be 0, in which case the default stream is used with associated synchronization rules.
|
||||
* @param [in] kernelParams pointer to kernel parameters.
|
||||
* @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and
|
||||
@@ -53,11 +53,11 @@ THE SOFTWARE.
|
||||
* the kernel launch. The event must be created before calling this API.
|
||||
* @param [in] stopEvent If non-null, specified event will be updated to track the stop time of
|
||||
* the kernel launch. The event must be created before calling this API.
|
||||
* @param [in] flags. The value of hipExtAnyOrderLaunch, signifies if kernel can be
|
||||
* @param [in] flags The value of hipExtAnyOrderLaunch, signifies if kernel can be
|
||||
* launched in any order.
|
||||
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue.
|
||||
* @returns #hipSuccess, #hipInvalidDeviceId, #hipErrorNotInitialized, #hipErrorInvalidValue.
|
||||
*
|
||||
* @warning kernellParams argument is not yet implemented in HIP, use extra instead.
|
||||
* @warning kernellParams argument is not yet implemented in HIP, use extra instead.
|
||||
* Please refer to hip_porting_driver_api.md for sample usage.
|
||||
* HIP/ROCm actually updates the start event when the associated kernel completes.
|
||||
* Currently, timing between startEvent and stopEvent does not include the time it takes to perform
|
||||
@@ -104,9 +104,9 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
|
||||
* the kernel launch. The event must be created before calling this API.
|
||||
* @param [in] stopEvent If non-null, specified event will be updated to track the stop time of
|
||||
* the kernel launch. The event must be created before calling this API.
|
||||
* @param [in] flags. The value of hipExtAnyOrderLaunch, signifies if kernel can be
|
||||
* @param [in] flags The value of hipExtAnyOrderLaunch, signifies if kernel can be
|
||||
* launched in any order.
|
||||
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue.
|
||||
* @returns #hipSuccess, #hipInvalidDeviceId, #hipErrorNotInitialized, #hipErrorInvalidValue.
|
||||
*
|
||||
*/
|
||||
extern "C" hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks,
|
||||
@@ -117,9 +117,9 @@ extern "C" hipError_t hipExtLaunchKernel(const void* function_address, dim3 numB
|
||||
/**
|
||||
* @brief Launches kernel with dimention parameters and shared memory on stream with templated kernel and arguments.
|
||||
*
|
||||
* @param [in] f Kernel to launch.
|
||||
* @param [in] numBlocks const number of blocks.
|
||||
* @param [in] dimBlocks const dimension of a block.
|
||||
* @param [in] kernel Kernel to launch.
|
||||
* @param [in] numBlocks const number of blocks.
|
||||
* @param [in] dimBlocks const dimension of a block.
|
||||
* @param [in] sharedMemBytes Amount of dynamic shared memory to allocate for this kernel.
|
||||
* HIP-Clang compiler provides support for extern shared declarations.
|
||||
* @param [in] stream Stream where the kernel should be dispatched.
|
||||
@@ -128,10 +128,11 @@ extern "C" hipError_t hipExtLaunchKernel(const void* function_address, dim3 numB
|
||||
* the kernel launch. The event must be created before calling this API.
|
||||
* @param [in] stopEvent If non-null, specified event will be updated to track the stop time of
|
||||
* the kernel launch. The event must be created before calling this API.
|
||||
* @param [in] flags. The value of hipExtAnyOrderLaunch, signifies if kernel can be
|
||||
* @param [in] flags The value of hipExtAnyOrderLaunch, signifies if kernel can be
|
||||
* launched in any order.
|
||||
* @param [in] args templated kernel arguments.
|
||||
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue.
|
||||
* @param [in] args templated kernel arguments.
|
||||
*
|
||||
*
|
||||
* Please refer to the application for sample usage at,
|
||||
* (https://github.com/ROCm-Developer-Tools/HIP/blob/rocm-4.5.x/tests/src/kernel/hipExtLaunchKernelGGL.cpp).
|
||||
*/
|
||||
|
||||
Rozdílový obsah nebyl zobrazen, protože je příliš veliký
Načíst rozdílové porovnání
@@ -29,11 +29,6 @@ THE SOFTWARE.
|
||||
#elif (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && \
|
||||
!(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
|
||||
|
||||
/**
|
||||
* @addtogroup Runtime Runtime Compilation
|
||||
* @{
|
||||
* @ingroup Runtime
|
||||
*/
|
||||
|
||||
#ifdef __cplusplus
|
||||
extern "C" {
|
||||
@@ -46,86 +41,93 @@ extern "C" {
|
||||
#endif
|
||||
|
||||
/**
|
||||
* @brief hiprtcResult
|
||||
* @enum
|
||||
*
|
||||
* @addtogroup GlobalDefs
|
||||
* @{
|
||||
*
|
||||
*/
|
||||
/**
|
||||
* hiprtc error code
|
||||
*/
|
||||
|
||||
typedef enum hiprtcResult {
|
||||
HIPRTC_SUCCESS = 0,
|
||||
HIPRTC_ERROR_OUT_OF_MEMORY = 1,
|
||||
HIPRTC_ERROR_PROGRAM_CREATION_FAILURE = 2,
|
||||
HIPRTC_ERROR_INVALID_INPUT = 3,
|
||||
HIPRTC_ERROR_INVALID_PROGRAM = 4,
|
||||
HIPRTC_ERROR_INVALID_OPTION = 5,
|
||||
HIPRTC_ERROR_COMPILATION = 6,
|
||||
HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7,
|
||||
HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8,
|
||||
HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9,
|
||||
HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10,
|
||||
HIPRTC_ERROR_INTERNAL_ERROR = 11,
|
||||
HIPRTC_ERROR_LINKING = 100
|
||||
HIPRTC_SUCCESS = 0, ///< Success
|
||||
HIPRTC_ERROR_OUT_OF_MEMORY = 1, ///< Out of memory
|
||||
HIPRTC_ERROR_PROGRAM_CREATION_FAILURE = 2, ///< Failed to create program
|
||||
HIPRTC_ERROR_INVALID_INPUT = 3, ///< Invalid input
|
||||
HIPRTC_ERROR_INVALID_PROGRAM = 4, ///< Invalid program
|
||||
HIPRTC_ERROR_INVALID_OPTION = 5, ///< Invalid option
|
||||
HIPRTC_ERROR_COMPILATION = 6, ///< Compilation error
|
||||
HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7, ///< Failed in builtin operation
|
||||
HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8, ///< No name expression after compilation
|
||||
HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9, ///< No lowered names before compilation
|
||||
HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10, ///< Invalid name expression
|
||||
HIPRTC_ERROR_INTERNAL_ERROR = 11, ///< Internal error
|
||||
HIPRTC_ERROR_LINKING = 100 ///< Error in linking
|
||||
} hiprtcResult;
|
||||
|
||||
/**
|
||||
* @brief hiprtcJIT_option
|
||||
* @enum
|
||||
*
|
||||
* hiprtc JIT option
|
||||
*/
|
||||
|
||||
typedef enum hiprtcJIT_option {
|
||||
HIPRTC_JIT_MAX_REGISTERS = 0,
|
||||
HIPRTC_JIT_THREADS_PER_BLOCK,
|
||||
HIPRTC_JIT_WALL_TIME,
|
||||
HIPRTC_JIT_INFO_LOG_BUFFER,
|
||||
HIPRTC_JIT_INFO_LOG_BUFFER_SIZE_BYTES,
|
||||
HIPRTC_JIT_ERROR_LOG_BUFFER,
|
||||
HIPRTC_JIT_ERROR_LOG_BUFFER_SIZE_BYTES,
|
||||
HIPRTC_JIT_OPTIMIZATION_LEVEL,
|
||||
HIPRTC_JIT_TARGET_FROM_HIPCONTEXT,
|
||||
HIPRTC_JIT_TARGET,
|
||||
HIPRTC_JIT_FALLBACK_STRATEGY,
|
||||
HIPRTC_JIT_GENERATE_DEBUG_INFO,
|
||||
HIPRTC_JIT_LOG_VERBOSE,
|
||||
HIPRTC_JIT_GENERATE_LINE_INFO,
|
||||
HIPRTC_JIT_CACHE_MODE,
|
||||
HIPRTC_JIT_NEW_SM3X_OPT,
|
||||
HIPRTC_JIT_FAST_COMPILE,
|
||||
HIPRTC_JIT_GLOBAL_SYMBOL_NAMES,
|
||||
HIPRTC_JIT_GLOBAL_SYMBOL_ADDRESS,
|
||||
HIPRTC_JIT_GLOBAL_SYMBOL_COUNT,
|
||||
HIPRTC_JIT_LTO,
|
||||
HIPRTC_JIT_FTZ,
|
||||
HIPRTC_JIT_PREC_DIV,
|
||||
HIPRTC_JIT_PREC_SQRT,
|
||||
HIPRTC_JIT_FMA,
|
||||
HIPRTC_JIT_NUM_OPTIONS,
|
||||
HIPRTC_JIT_IR_TO_ISA_OPT_EXT = 10000, //! AMD only. Linker options to be passed on to
|
||||
HIPRTC_JIT_IR_TO_ISA_OPT_COUNT_EXT, //! AMD only. Count of linker options
|
||||
HIPRTC_JIT_MAX_REGISTERS = 0, ///< Maximum registers
|
||||
HIPRTC_JIT_THREADS_PER_BLOCK, ///< Thread per block
|
||||
HIPRTC_JIT_WALL_TIME, ///< Time from aall clock
|
||||
HIPRTC_JIT_INFO_LOG_BUFFER, ///< Log buffer info
|
||||
HIPRTC_JIT_INFO_LOG_BUFFER_SIZE_BYTES, ///< Log buffer size in bytes
|
||||
HIPRTC_JIT_ERROR_LOG_BUFFER, ///< Log buffer error
|
||||
HIPRTC_JIT_ERROR_LOG_BUFFER_SIZE_BYTES, ///< Log buffer size in bytes
|
||||
HIPRTC_JIT_OPTIMIZATION_LEVEL, ///< Optimization level
|
||||
HIPRTC_JIT_TARGET_FROM_HIPCONTEXT, ///<
|
||||
HIPRTC_JIT_TARGET, ///< JIT target
|
||||
HIPRTC_JIT_FALLBACK_STRATEGY, ///< Fallback strategy
|
||||
HIPRTC_JIT_GENERATE_DEBUG_INFO, ///< Generate debug information
|
||||
HIPRTC_JIT_LOG_VERBOSE, ///< Log verbose
|
||||
HIPRTC_JIT_GENERATE_LINE_INFO, ///< Generate line information
|
||||
HIPRTC_JIT_CACHE_MODE, ///< Cache mode
|
||||
HIPRTC_JIT_NEW_SM3X_OPT, ///< New SM3X option
|
||||
HIPRTC_JIT_FAST_COMPILE, ///< Fast compile
|
||||
HIPRTC_JIT_GLOBAL_SYMBOL_NAMES, ///< Global symbol names
|
||||
HIPRTC_JIT_GLOBAL_SYMBOL_ADDRESS, ///< Global symbol address
|
||||
HIPRTC_JIT_GLOBAL_SYMBOL_COUNT, ///< Global symbol count
|
||||
HIPRTC_JIT_LTO, ///< LTO
|
||||
HIPRTC_JIT_FTZ, ///< FTZ
|
||||
HIPRTC_JIT_PREC_DIV, ///< Prec_VIV
|
||||
HIPRTC_JIT_PREC_SQRT, ///< PREC_SQRT
|
||||
HIPRTC_JIT_FMA, ///< FMA
|
||||
HIPRTC_JIT_NUM_OPTIONS, ///< Number of options
|
||||
HIPRTC_JIT_IR_TO_ISA_OPT_EXT = 10000, //< AMD only. Linker options to be passed on to
|
||||
HIPRTC_JIT_IR_TO_ISA_OPT_COUNT_EXT, //< AMD only. Count of linker options
|
||||
} hiprtcJIT_option;
|
||||
|
||||
/**
|
||||
* @brief hiprtcJITInputType
|
||||
* @enum
|
||||
*
|
||||
* hiprtc JIT input type
|
||||
*/
|
||||
typedef enum hiprtcJITInputType {
|
||||
HIPRTC_JIT_INPUT_CUBIN = 0,
|
||||
HIPRTC_JIT_INPUT_PTX,
|
||||
HIPRTC_JIT_INPUT_FATBINARY,
|
||||
HIPRTC_JIT_INPUT_OBJECT,
|
||||
HIPRTC_JIT_INPUT_LIBRARY,
|
||||
HIPRTC_JIT_INPUT_NVVM,
|
||||
HIPRTC_JIT_NUM_LEGACY_INPUT_TYPES,
|
||||
HIPRTC_JIT_INPUT_LLVM_BITCODE = 100,
|
||||
HIPRTC_JIT_INPUT_LLVM_BUNDLED_BITCODE = 101,
|
||||
HIPRTC_JIT_INPUT_LLVM_ARCHIVES_OF_BUNDLED_BITCODE = 102,
|
||||
HIPRTC_JIT_INPUT_CUBIN = 0, ///< Input cubin
|
||||
HIPRTC_JIT_INPUT_PTX, ///< Input PTX
|
||||
HIPRTC_JIT_INPUT_FATBINARY, ///< Input fat binary
|
||||
HIPRTC_JIT_INPUT_OBJECT, ///< Input object
|
||||
HIPRTC_JIT_INPUT_LIBRARY, ///< Input library
|
||||
HIPRTC_JIT_INPUT_NVVM, ///< Input NVVM
|
||||
HIPRTC_JIT_NUM_LEGACY_INPUT_TYPES, ///< Number of legacy input type
|
||||
HIPRTC_JIT_INPUT_LLVM_BITCODE = 100, ///< LLVM bitcode
|
||||
HIPRTC_JIT_INPUT_LLVM_BUNDLED_BITCODE = 101, ///< LLVM bundled bitcode
|
||||
HIPRTC_JIT_INPUT_LLVM_ARCHIVES_OF_BUNDLED_BITCODE = 102, ///< LLVM archives of boundled bitcode
|
||||
HIPRTC_JIT_NUM_INPUT_TYPES = (HIPRTC_JIT_NUM_LEGACY_INPUT_TYPES + 3)
|
||||
} hiprtcJITInputType;
|
||||
|
||||
typedef struct ihiprtcLinkState* hiprtcLinkState;
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
/**
|
||||
* hiprtc link state
|
||||
*
|
||||
*/
|
||||
typedef struct ihiprtcLinkState* hiprtcLinkState;
|
||||
/**
|
||||
* @ingroup Runtime
|
||||
*
|
||||
* @brief Returns text string message to explain the error which occurred
|
||||
*
|
||||
* @param [in] result code to convert to string.
|
||||
@@ -139,6 +141,7 @@ typedef struct ihiprtcLinkState* hiprtcLinkState;
|
||||
const char* hiprtcGetErrorString(hiprtcResult result);
|
||||
|
||||
/**
|
||||
* @ingroup Runtime
|
||||
* @brief Sets the parameters as major and minor version.
|
||||
*
|
||||
* @param [out] major HIP Runtime Compilation major version.
|
||||
@@ -146,38 +149,47 @@ const char* hiprtcGetErrorString(hiprtcResult result);
|
||||
*
|
||||
*/
|
||||
hiprtcResult hiprtcVersion(int* major, int* minor);
|
||||
|
||||
/**
|
||||
* hiprtc program
|
||||
*
|
||||
*/
|
||||
typedef struct _hiprtcProgram* hiprtcProgram;
|
||||
|
||||
/**
|
||||
* @ingroup Runtime
|
||||
* @brief Adds the given name exprssion to the runtime compilation program.
|
||||
*
|
||||
* @param [in] prog runtime compilation program instance.
|
||||
* @param [in] name_expression const char pointer to the name expression.
|
||||
* @return HIPRTC_SUCCESS
|
||||
* @return #HIPRTC_SUCCESS
|
||||
*
|
||||
* If const char pointer is NULL, it will return HIPRTC_ERROR_INVALID_INPUT.
|
||||
* If const char pointer is NULL, it will return #HIPRTC_ERROR_INVALID_INPUT.
|
||||
*
|
||||
* @see hiprtcResult
|
||||
*/
|
||||
hiprtcResult hiprtcAddNameExpression(hiprtcProgram prog, const char* name_expression);
|
||||
hiprtcResult hiprtcAddNameExpression(hiprtcProgram prog,
|
||||
const char* name_expression);
|
||||
|
||||
/**
|
||||
* @ingroup Runtime
|
||||
* @brief Compiles the given runtime compilation program.
|
||||
*
|
||||
* @param [in] prog runtime compilation program instance.
|
||||
* @param [in] numOptions number of compiler options.
|
||||
* @param [in] options compiler options as const array of strins.
|
||||
* @return HIPRTC_SUCCESS
|
||||
* @return #HIPRTC_SUCCESS
|
||||
*
|
||||
* If the compiler failed to build the runtime compilation program,
|
||||
* it will return HIPRTC_ERROR_COMPILATION.
|
||||
* it will return #HIPRTC_ERROR_COMPILATION.
|
||||
*
|
||||
* @see hiprtcResult
|
||||
*/
|
||||
hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, int numOptions, const char** options);
|
||||
hiprtcResult hiprtcCompileProgram(hiprtcProgram prog,
|
||||
int numOptions,
|
||||
const char** options);
|
||||
|
||||
/**
|
||||
* @ingroup Runtime
|
||||
* @brief Creates an instance of hiprtcProgram with the given input parameters,
|
||||
* and sets the output hiprtcProgram prog with it.
|
||||
*
|
||||
@@ -187,53 +199,58 @@ hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, int numOptions, const char
|
||||
* @param [in] numHeaders number of headers.
|
||||
* @param [in] headers array of strings pointing to headers.
|
||||
* @param [in] includeNames array of strings pointing to names included in program source.
|
||||
* @return HIPRTC_SUCCESS
|
||||
* @return #HIPRTC_SUCCESS
|
||||
*
|
||||
* Any invalide input parameter, it will return HIPRTC_ERROR_INVALID_INPUT
|
||||
* or HIPRTC_ERROR_INVALID_PROGRAM.
|
||||
* Any invalide input parameter, it will return #HIPRTC_ERROR_INVALID_INPUT
|
||||
* or #HIPRTC_ERROR_INVALID_PROGRAM.
|
||||
*
|
||||
* If failed to create the program, it will return HIPRTC_ERROR_PROGRAM_CREATION_FAILURE.
|
||||
* If failed to create the program, it will return #HIPRTC_ERROR_PROGRAM_CREATION_FAILURE.
|
||||
*
|
||||
* @see hiprtcResult
|
||||
*/
|
||||
hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, const char* src, const char* name,
|
||||
int numHeaders, const char** headers, const char** includeNames);
|
||||
hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog,
|
||||
const char* src,
|
||||
const char* name,
|
||||
int numHeaders,
|
||||
const char** headers,
|
||||
const char** includeNames);
|
||||
|
||||
/**
|
||||
* @brief Destroys an instance of given hiprtcProgram.
|
||||
*
|
||||
* @ingroup Runtime
|
||||
* @param [in] prog runtime compilation program instance.
|
||||
* @return HIPRTC_SUCCESS
|
||||
* @return #HIPRTC_SUCCESS
|
||||
*
|
||||
* If prog is NULL, it will return HIPRTC_ERROR_INVALID_INPUT.
|
||||
* If prog is NULL, it will return #HIPRTC_ERROR_INVALID_INPUT.
|
||||
*
|
||||
* @see hiprtcResult
|
||||
*/
|
||||
hiprtcResult hiprtcDestroyProgram(hiprtcProgram* prog);
|
||||
|
||||
/**
|
||||
* @brief Gets the lowered (mangled) name from an instance of hiprtcProgram with the given input
|
||||
* parameters, and sets the output lowered_name with it.
|
||||
*
|
||||
* @brief Gets the lowered (mangled) name from an instance of hiprtcProgram with the given input parameters,
|
||||
* and sets the output lowered_name with it.
|
||||
* @ingroup Runtime
|
||||
* @param [in] prog runtime compilation program instance.
|
||||
* @param [in] name_expression const char pointer to the name expression.
|
||||
* @param [in, out] lowered_name const char array to the lowered (mangled) name.
|
||||
* @return HIPRTC_SUCCESS
|
||||
* @return #HIPRTC_SUCCESS
|
||||
*
|
||||
* If any invalide nullptr input parameters, it will return HIPRTC_ERROR_INVALID_INPUT
|
||||
* If any invalide nullptr input parameters, it will return #HIPRTC_ERROR_INVALID_INPUT
|
||||
*
|
||||
* If name_expression is not found, it will return HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID
|
||||
* If name_expression is not found, it will return #HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID
|
||||
*
|
||||
* If failed to get lowered_name from the program, it will return HIPRTC_ERROR_COMPILATION.
|
||||
* If failed to get lowered_name from the program, it will return #HIPRTC_ERROR_COMPILATION.
|
||||
*
|
||||
* @see hiprtcResult
|
||||
*/
|
||||
hiprtcResult hiprtcGetLoweredName(hiprtcProgram prog, const char* name_expression,
|
||||
hiprtcResult hiprtcGetLoweredName(hiprtcProgram prog,
|
||||
const char* name_expression,
|
||||
const char** lowered_name);
|
||||
|
||||
/**
|
||||
* @brief Gets the log generated by the runtime compilation program instance.
|
||||
*
|
||||
* @ingroup Runtime
|
||||
* @param [in] prog runtime compilation program instance.
|
||||
* @param [out] log memory pointer to the generated log.
|
||||
* @return HIPRTC_SUCCESS
|
||||
@@ -251,11 +268,12 @@ hiprtcResult hiprtcGetProgramLog(hiprtcProgram prog, char* log);
|
||||
*
|
||||
* @see hiprtcResult
|
||||
*/
|
||||
hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram prog, size_t* logSizeRet);
|
||||
hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram prog,
|
||||
size_t* logSizeRet);
|
||||
|
||||
/**
|
||||
* @brief Gets the pointer of compilation binary by the runtime compilation program instance.
|
||||
*
|
||||
* @ingroup Runtime
|
||||
* @param [in] prog runtime compilation program instance.
|
||||
* @param [out] code char pointer to binary.
|
||||
* @return HIPRTC_SUCCESS
|
||||
@@ -266,9 +284,9 @@ hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code);
|
||||
|
||||
/**
|
||||
* @brief Gets the size of compilation binary by the runtime compilation program instance.
|
||||
*
|
||||
* @ingroup Runtime
|
||||
* @param [in] prog runtime compilation program instance.
|
||||
* @param [out] code the size of binary.
|
||||
* @param [out] codeSizeRet the size of binary.
|
||||
* @return HIPRTC_SUCCESS
|
||||
*
|
||||
* @see hiprtcResult
|
||||
@@ -279,7 +297,7 @@ hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, size_t* codeSizeRet);
|
||||
* @brief Gets the pointer of compiled bitcode by the runtime compilation program instance.
|
||||
*
|
||||
* @param [in] prog runtime compilation program instance.
|
||||
* @param [out] code char pointer to bitcode.
|
||||
* @param [out] bitcode char pointer to bitcode.
|
||||
* @return HIPRTC_SUCCESS
|
||||
*
|
||||
* @see hiprtcResult
|
||||
@@ -288,11 +306,11 @@ hiprtcResult hiprtcGetBitcode(hiprtcProgram prog, char* bitcode);
|
||||
|
||||
/**
|
||||
* @brief Gets the size of compiled bitcode by the runtime compilation program instance.
|
||||
*
|
||||
* @ingroup Runtime
|
||||
*
|
||||
* @param [in] prog runtime compilation program instance.
|
||||
* @param [out] code the size of bitcode.
|
||||
* @return HIPRTC_SUCCESS
|
||||
* @param [out] bitcode_size the size of bitcode.
|
||||
* @return #HIPRTC_SUCCESS
|
||||
*
|
||||
* @see hiprtcResult
|
||||
*/
|
||||
@@ -300,10 +318,13 @@ hiprtcResult hiprtcGetBitcodeSize(hiprtcProgram prog, size_t* bitcode_size);
|
||||
|
||||
/**
|
||||
* @brief Creates the link instance via hiprtc APIs.
|
||||
*
|
||||
* @param [in] hip_jit_options
|
||||
* @param [out] hiprtc link state instance
|
||||
* @return HIPRTC_SUCCESS
|
||||
* @ingroup Runtime
|
||||
* @param [in] num_options Number of options
|
||||
* @param [in] option_ptr Array of options
|
||||
* @param [in] option_vals_pptr Array of option values cast to void*
|
||||
* @param [out] hip_link_state_ptr hiprtc link state created upon success
|
||||
*
|
||||
* @return #HIPRTC_SUCCESS, #HIPRTC_ERROR_INVALID_INPUT, #HIPRTC_ERROR_INVALID_OPTION
|
||||
*
|
||||
* @see hiprtcResult
|
||||
*/
|
||||
@@ -312,14 +333,18 @@ hiprtcResult hiprtcLinkCreate(unsigned int num_options, hiprtcJIT_option* option
|
||||
|
||||
/**
|
||||
* @brief Adds a file with bit code to be linked with options
|
||||
* @ingroup Runtime
|
||||
* @param [in] hip_link_state hiprtc link state
|
||||
* @param [in] input_type Type of the input data or bitcode
|
||||
* @param [in] file_path Path to the input file where bitcode is present
|
||||
* @param [in] num_options Size of the options
|
||||
* @param [in] options_ptr Array of options applied to this input
|
||||
* @param [in] option_values Array of option values cast to void*
|
||||
*
|
||||
* @param [in] hiprtc link state, jit input type, file path,
|
||||
* option reated parameters.
|
||||
* @param [out] None.
|
||||
* @return HIPRTC_SUCCESS
|
||||
* @return #HIPRTC_SUCCESS
|
||||
*
|
||||
* If input values are invalid, it will
|
||||
* @return HIPRTC_ERROR_INVALID_INPUT
|
||||
* @return #HIPRTC_ERROR_INVALID_INPUT
|
||||
*
|
||||
* @see hiprtcResult
|
||||
*/
|
||||
@@ -330,14 +355,20 @@ hiprtcResult hiprtcLinkAddFile(hiprtcLinkState hip_link_state, hiprtcJITInputTyp
|
||||
|
||||
/**
|
||||
* @brief Completes the linking of the given program.
|
||||
* @ingroup Runtime
|
||||
* @param [in] hip_link_state hiprtc link state
|
||||
* @param [in] input_type Type of the input data or bitcode
|
||||
* @param [in] image Input data which is null terminated
|
||||
* @param [in] image_size Size of the input data
|
||||
* @param [in] name Optional name for this input
|
||||
* @param [in] num_options Size of the options
|
||||
* @param [in] options_ptr Array of options applied to this input
|
||||
* @param [in] option_values Array of option values cast to void*
|
||||
*
|
||||
* @param [in] hiprtc link state, jit input type, image_ptr ,
|
||||
* option reated parameters.
|
||||
* @param [out] None.
|
||||
* @return HIPRTC_SUCCESS
|
||||
* @return #HIPRTC_SUCCESS, #HIPRTC_ERROR_INVALID_INPUT
|
||||
*
|
||||
* If adding the file fails, it will
|
||||
* @return HIPRTC_ERROR_PROGRAM_CREATION_FAILURE
|
||||
* @return #HIPRTC_ERROR_PROGRAM_CREATION_FAILURE
|
||||
*
|
||||
* @see hiprtcResult
|
||||
*/
|
||||
@@ -349,13 +380,15 @@ hiprtcResult hiprtcLinkAddData(hiprtcLinkState hip_link_state, hiprtcJITInputTyp
|
||||
|
||||
/**
|
||||
* @brief Completes the linking of the given program.
|
||||
* @ingroup Runtime
|
||||
* @param [in] hip_link_state hiprtc link state
|
||||
* @param [out] bin_out Upon success, points to the output binary
|
||||
* @param [out] size_out Size of the binary is stored (optional)
|
||||
*
|
||||
* @param [in] hiprtc link state instance
|
||||
* @param [out] linked_binary, linked_binary_size.
|
||||
* @return HIPRTC_SUCCESS
|
||||
* @return #HIPRTC_SUCCESS
|
||||
*
|
||||
* If adding the data fails, it will
|
||||
* @return HIPRTC_ERROR_PROGRAM_CREATION_FAILURE
|
||||
* @return #HIPRTC_ERROR_LINKING
|
||||
*
|
||||
* @see hiprtcResult
|
||||
*/
|
||||
@@ -363,13 +396,10 @@ hiprtcResult hiprtcLinkComplete(hiprtcLinkState hip_link_state, void** bin_out,
|
||||
|
||||
/**
|
||||
* @brief Deletes the link instance via hiprtc APIs.
|
||||
* @ingroup Runtime
|
||||
* @param [in] hip_link_state link state instance
|
||||
*
|
||||
* @param [in] hiprtc link state instance
|
||||
* @param [out] code the size of binary.
|
||||
* @return HIPRTC_SUCCESS
|
||||
*
|
||||
* If linking fails, it will
|
||||
* @return HIPRTC_ERROR_LINKING
|
||||
* @return #HIPRTC_SUCCESS
|
||||
*
|
||||
* @see hiprtcResult
|
||||
*/
|
||||
@@ -382,10 +412,7 @@ hiprtcResult hiprtcLinkDestroy(hiprtcLinkState hip_link_state);
|
||||
#ifdef __cplusplus
|
||||
}
|
||||
#endif /* __cplusplus */
|
||||
// doxygen end HIPrtc feature
|
||||
/**
|
||||
* @}
|
||||
*/
|
||||
|
||||
#else
|
||||
#error("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
|
||||
#endif
|
||||
|
||||
@@ -1,52 +0,0 @@
|
||||
# Copyright (c) 2020 - 2021 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.
|
||||
|
||||
project(bit_extract)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if(NOT WIN32 AND NOT DEFINED __HIP_ENABLE_PCH)
|
||||
set(__HIP_ENABLE_PCH ON CACHE BOOL "enable/disable pre-compiled hip headers")
|
||||
endif()
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
if(${__HIP_ENABLE_PCH})
|
||||
add_definitions(-D__HIP_ENABLE_PCH)
|
||||
endif()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
# Set compiler and linker
|
||||
set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
|
||||
|
||||
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})
|
||||
|
||||
# Create the excutable
|
||||
add_executable(bit_extract bit_extract.cpp)
|
||||
|
||||
# Link with HIP
|
||||
target_link_libraries(bit_extract hip::host)
|
||||
@@ -1,46 +0,0 @@
|
||||
# Copyright (c) 2016 - 2021 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.
|
||||
|
||||
#Dependencies : [MYHIP]/bin must be in user's path.
|
||||
ifeq ($(OS),Windows_NT)
|
||||
$(error Makefile is not supported on windows platform. Please use cmake instead to build sample.)
|
||||
endif
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform)
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
|
||||
# Show how to use PLATFORM to specify different options for each compiler:
|
||||
ifeq (${HIP_PLATFORM}, nvcc)
|
||||
HIPCC_FLAGS = -gencode=arch=compute_20,code=sm_20
|
||||
endif
|
||||
|
||||
EXE=bit_extract
|
||||
|
||||
$(EXE): bit_extract.cpp
|
||||
$(HIPCC) $(HIPCC_FLAGS) $< -o $@
|
||||
|
||||
all: $(EXE)
|
||||
|
||||
clean:
|
||||
rm -f *.o $(EXE)
|
||||
@@ -1,6 +0,0 @@
|
||||
# bit_extract
|
||||
|
||||
Show an application written directly in HIP which uses platform-specific check on __HIP_PLATFORM_AMD__ to enable use of
|
||||
an instruction that only exists on the HCC platform.
|
||||
|
||||
See related [blog](http://gpuopen.com/platform-aware-coding-inside-hip/) demonstrating platform specialization.
|
||||
@@ -1,113 +0,0 @@
|
||||
/*
|
||||
Copyright (c) 2015 - 2021 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 <stdio.h>
|
||||
#include <iostream>
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
#define CHECK(cmd) \
|
||||
{ \
|
||||
hipError_t error = cmd; \
|
||||
if (error != hipSuccess) { \
|
||||
fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error, \
|
||||
__FILE__, __LINE__); \
|
||||
exit(EXIT_FAILURE); \
|
||||
} \
|
||||
}
|
||||
|
||||
__global__ void bit_extract_kernel(uint32_t* C_d, const uint32_t* A_d, size_t N) {
|
||||
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
|
||||
for (size_t i = offset; i < N; i += stride) {
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
C_d[i] = __bitextract_u32(A_d[i], 8, 4);
|
||||
#else /* defined __HIP_PLATFORM_NVIDIA__ or other path */
|
||||
C_d[i] = ((A_d[i] & 0xf00) >> 8);
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
uint32_t *A_d, *C_d;
|
||||
uint32_t *A_h, *C_h;
|
||||
size_t N = 1000000;
|
||||
size_t Nbytes = N * sizeof(uint32_t);
|
||||
|
||||
#ifdef __HIP_ENABLE_PCH
|
||||
// Verify hip_pch.o
|
||||
const char* pch = nullptr;
|
||||
unsigned int size = 0;
|
||||
__hipGetPCH(&pch, &size);
|
||||
printf("pch size: %u\n", size);
|
||||
if (size == 0) {
|
||||
printf("__hipGetPCH failed!\n");
|
||||
return -1;
|
||||
} else {
|
||||
printf("__hipGetPCH succeeded!\n");
|
||||
}
|
||||
#endif
|
||||
|
||||
int deviceId;
|
||||
CHECK(hipGetDevice(&deviceId));
|
||||
hipDeviceProp_t props;
|
||||
CHECK(hipGetDeviceProperties(&props, deviceId));
|
||||
printf("info: running on device #%d %s\n", deviceId, props.name);
|
||||
|
||||
|
||||
printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0);
|
||||
A_h = (uint32_t*)malloc(Nbytes);
|
||||
CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
C_h = (uint32_t*)malloc(Nbytes);
|
||||
CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
A_h[i] = i;
|
||||
}
|
||||
|
||||
printf("info: allocate device mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0);
|
||||
CHECK(hipMalloc(&A_d, Nbytes));
|
||||
CHECK(hipMalloc(&C_d, Nbytes));
|
||||
|
||||
printf("info: copy Host2Device\n");
|
||||
CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
printf("info: launch 'bit_extract_kernel' \n");
|
||||
const unsigned blocks = 512;
|
||||
const unsigned threadsPerBlock = 256;
|
||||
hipLaunchKernelGGL(bit_extract_kernel, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
|
||||
|
||||
printf("info: copy Device2Host\n");
|
||||
CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
printf("info: check result\n");
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
unsigned Agold = ((A_h[i] & 0xf00) >> 8);
|
||||
if (C_h[i] != Agold) {
|
||||
fprintf(stderr, "mismatch detected.\n");
|
||||
printf("%zu: %08x =? %08x (Ain=%08x)\n", i, C_h[i], Agold, A_h[i]);
|
||||
CHECK(hipErrorUnknown);
|
||||
}
|
||||
}
|
||||
printf("PASSED!\n");
|
||||
}
|
||||
@@ -1,60 +0,0 @@
|
||||
# Copyright (c) 2020 - 2021 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.
|
||||
|
||||
project(module_api)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
# Set compiler and linker
|
||||
set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
|
||||
|
||||
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})
|
||||
|
||||
# Create the excutable
|
||||
add_executable(runKernel.hip.out runKernel.cpp)
|
||||
add_executable(launchKernelHcc.hip.out launchKernelHcc.cpp)
|
||||
add_executable(defaultDriver.hip.out defaultDriver.cpp)
|
||||
|
||||
# Generate code object
|
||||
add_custom_target(
|
||||
codeobj
|
||||
ALL
|
||||
COMMAND ${HIP_HIPCC_EXECUTABLE} --genco ../vcpy_kernel.cpp -o vcpy_kernel.code
|
||||
COMMENT "codeobj generated"
|
||||
)
|
||||
|
||||
add_dependencies(runKernel.hip.out codeobj)
|
||||
add_dependencies(launchKernelHcc.hip.out codeobj)
|
||||
add_dependencies(defaultDriver.hip.out codeobj)
|
||||
|
||||
# Link with HIP
|
||||
target_link_libraries(runKernel.hip.out hip::host)
|
||||
target_link_libraries(launchKernelHcc.hip.out hip::host)
|
||||
target_link_libraries(defaultDriver.hip.out hip::host)
|
||||
@@ -1,46 +0,0 @@
|
||||
# Copyright (c) 2016 - 2021 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.
|
||||
ifeq ($(OS),Windows_NT)
|
||||
$(error Makefile is not supported on windows platform. Please use cmake instead to build sample.)
|
||||
endif
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
|
||||
|
||||
all: vcpy_kernel.code runKernel.hip.out launchKernelHcc.hip.out defaultDriver.hip.out
|
||||
|
||||
runKernel.hip.out: runKernel.cpp
|
||||
$(HIPCC) $(HIPCC_FLAGS) $< -o $@
|
||||
|
||||
launchKernelHcc.hip.out: launchKernelHcc.cpp
|
||||
$(HIPCC) $(HIPCC_FLAGS) $< -o $@
|
||||
|
||||
defaultDriver.hip.out: defaultDriver.cpp
|
||||
$(HIPCC) $(HIPCC_FLAGS) $< -o $@
|
||||
|
||||
vcpy_kernel.code: vcpy_kernel.cpp
|
||||
$(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@
|
||||
|
||||
clean:
|
||||
rm -f *.code *.out
|
||||
@@ -1,87 +0,0 @@
|
||||
/*
|
||||
Copyright (c) 2015 - 2021 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 "hip/hip_runtime.h"
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <vector>
|
||||
|
||||
#define LEN 64
|
||||
#define SIZE LEN << 2
|
||||
|
||||
#define fileName "vcpy_kernel.code"
|
||||
#define kernel_name "hello_world"
|
||||
|
||||
int main() {
|
||||
float *A, *B;
|
||||
hipDeviceptr_t Ad, Bd;
|
||||
A = new float[LEN];
|
||||
B = new float[LEN];
|
||||
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
A[i] = i * 1.0f;
|
||||
B[i] = 0.0f;
|
||||
}
|
||||
|
||||
hipInit(0);
|
||||
hipDevice_t device;
|
||||
hipCtx_t context;
|
||||
hipDeviceGet(&device, 0);
|
||||
hipCtxCreate(&context, 0, device);
|
||||
|
||||
hipMalloc((void**)&Ad, SIZE);
|
||||
hipMalloc((void**)&Bd, SIZE);
|
||||
|
||||
hipMemcpyHtoD(Ad, A, SIZE);
|
||||
hipMemcpyHtoD(Bd, B, SIZE);
|
||||
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
hipModuleLoad(&Module, fileName);
|
||||
hipModuleGetFunction(&Function, Module, kernel_name);
|
||||
|
||||
void* args[2] = {&Ad, &Bd};
|
||||
|
||||
hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, args, nullptr);
|
||||
|
||||
hipMemcpyDtoH(B, Bd, SIZE);
|
||||
int mismatchCount = 0;
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
if (A[i] != B[i]) {
|
||||
mismatchCount++;
|
||||
std::cout << "error: mismatch " << A[i] << " != " << B[i] << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
if (mismatchCount == 0) {
|
||||
std::cout << "PASSED!\n";
|
||||
} else {
|
||||
std::cout << "FAILED!\n";
|
||||
};
|
||||
|
||||
hipFree(Ad);
|
||||
hipFree(Bd);
|
||||
delete[] A;
|
||||
delete[] B;
|
||||
hipCtxDestroy(context);
|
||||
return 0;
|
||||
}
|
||||
@@ -1,111 +0,0 @@
|
||||
/*
|
||||
Copyright (c) 2015 - 2021 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 "hip/hip_runtime.h"
|
||||
#include "hip/hip_runtime_api.h"
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <vector>
|
||||
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
#include <hip/hip_ext.h>
|
||||
#endif
|
||||
|
||||
#define LEN 64
|
||||
#define SIZE LEN << 2
|
||||
|
||||
#define fileName "vcpy_kernel.code"
|
||||
#define kernel_name "hello_world"
|
||||
|
||||
#define HIP_CHECK(status) \
|
||||
if (status != hipSuccess) { \
|
||||
std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl; \
|
||||
exit(0); \
|
||||
}
|
||||
|
||||
int main() {
|
||||
float *A, *B;
|
||||
hipDeviceptr_t Ad, Bd;
|
||||
A = new float[LEN];
|
||||
B = new float[LEN];
|
||||
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
A[i] = i * 1.0f;
|
||||
B[i] = 0.0f;
|
||||
}
|
||||
|
||||
hipInit(0);
|
||||
hipDevice_t device;
|
||||
hipCtx_t context;
|
||||
hipDeviceGet(&device, 0);
|
||||
hipCtxCreate(&context, 0, device);
|
||||
|
||||
hipMalloc((void**)&Ad, SIZE);
|
||||
hipMalloc((void**)&Bd, SIZE);
|
||||
|
||||
hipMemcpyHtoD(Ad, A, SIZE);
|
||||
hipMemcpyHtoD(Bd, B, SIZE);
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
HIP_CHECK(hipModuleLoad(&Module, fileName));
|
||||
HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
|
||||
|
||||
struct {
|
||||
void* _Ad;
|
||||
void* _Bd;
|
||||
} args;
|
||||
|
||||
args._Ad = Ad;
|
||||
args._Bd = Bd;
|
||||
|
||||
|
||||
size_t size = sizeof(args);
|
||||
|
||||
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END};
|
||||
|
||||
HIP_CHECK(
|
||||
hipExtModuleLaunchKernel(Function, LEN, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config, 0));
|
||||
|
||||
hipMemcpyDtoH(B, Bd, SIZE);
|
||||
|
||||
int mismatchCount = 0;
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
if (A[i] != B[i]) {
|
||||
mismatchCount++;
|
||||
std::cout << "error: mismatch " << A[i] << " != " << B[i] << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
if (mismatchCount == 0) {
|
||||
std::cout << "PASSED!\n";
|
||||
} else {
|
||||
std::cout << "FAILED!\n";
|
||||
};
|
||||
|
||||
hipFree(Ad);
|
||||
hipFree(Bd);
|
||||
delete[] A;
|
||||
delete[] B;
|
||||
hipCtxDestroy(context);
|
||||
return 0;
|
||||
}
|
||||
@@ -1,106 +0,0 @@
|
||||
/*
|
||||
Copyright (c) 2015 - 2021 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 "hip/hip_runtime.h"
|
||||
#include "hip/hip_runtime_api.h"
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <vector>
|
||||
#include <hip/hip_hcc.h>
|
||||
|
||||
#define LEN 64
|
||||
#define SIZE LEN << 2
|
||||
|
||||
#define fileName "vcpy_kernel.code"
|
||||
#define kernel_name "hello_world"
|
||||
|
||||
#define HIP_CHECK(status) \
|
||||
if (status != hipSuccess) { \
|
||||
std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl; \
|
||||
exit(0); \
|
||||
}
|
||||
|
||||
int main() {
|
||||
float *A, *B;
|
||||
hipDeviceptr_t Ad, Bd;
|
||||
A = new float[LEN];
|
||||
B = new float[LEN];
|
||||
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
A[i] = i * 1.0f;
|
||||
B[i] = 0.0f;
|
||||
}
|
||||
|
||||
hipInit(0);
|
||||
hipDevice_t device;
|
||||
hipCtx_t context;
|
||||
hipDeviceGet(&device, 0);
|
||||
hipCtxCreate(&context, 0, device);
|
||||
|
||||
hipMalloc((void**)&Ad, SIZE);
|
||||
hipMalloc((void**)&Bd, SIZE);
|
||||
|
||||
hipMemcpyHtoD(Ad, A, SIZE);
|
||||
hipMemcpyHtoD(Bd, B, SIZE);
|
||||
hipModule_t Module;
|
||||
hipFunction_t Function;
|
||||
HIP_CHECK(hipModuleLoad(&Module, fileName));
|
||||
HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
|
||||
|
||||
struct {
|
||||
void* _Ad;
|
||||
void* _Bd;
|
||||
} args;
|
||||
|
||||
args._Ad = (void*) Ad;
|
||||
args._Bd = (void*) Bd;
|
||||
|
||||
size_t size = sizeof(args);
|
||||
|
||||
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END};
|
||||
|
||||
HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));
|
||||
|
||||
hipMemcpyDtoH(B, Bd, SIZE);
|
||||
|
||||
int mismatchCount = 0;
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
if (A[i] != B[i]) {
|
||||
mismatchCount++;
|
||||
std::cout << "error: mismatch " << A[i] << " != " << B[i] << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
if (mismatchCount == 0) {
|
||||
std::cout << "PASSED!\n";
|
||||
} else {
|
||||
std::cout << "FAILED!\n";
|
||||
};
|
||||
|
||||
hipFree(Ad);
|
||||
hipFree(Bd);
|
||||
delete[] A;
|
||||
delete[] B;
|
||||
hipCtxDestroy(context);
|
||||
return 0;
|
||||
}
|
||||
@@ -1,28 +0,0 @@
|
||||
/*
|
||||
Copyright (c) 2015 - 2021 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 "hip/hip_runtime.h"
|
||||
|
||||
extern "C" __global__ void hello_world(float* a, float* b) {
|
||||
int tx = threadIdx.x;
|
||||
b[tx] = a[tx];
|
||||
}
|
||||
@@ -1,54 +0,0 @@
|
||||
# Copyright (c) 2020 - 2021 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.
|
||||
|
||||
project(modile_api_global)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
# Set compiler and linker
|
||||
set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
|
||||
|
||||
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})
|
||||
|
||||
# Create the excutable
|
||||
add_executable(runKernel.hip.out runKernel.cpp)
|
||||
|
||||
# Generate code object
|
||||
add_custom_target(
|
||||
codeobj
|
||||
ALL
|
||||
COMMAND ${HIP_HIPCC_EXECUTABLE} --genco ../vcpy_kernel.cpp -o vcpy_kernel.code
|
||||
COMMENT "codeobj generated"
|
||||
)
|
||||
|
||||
add_dependencies(runKernel.hip.out codeobj)
|
||||
|
||||
# Link with HIP
|
||||
target_link_libraries(runKernel.hip.out hip::host)
|
||||
@@ -1,40 +0,0 @@
|
||||
# Copyright (c) 2017 - 2021 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.
|
||||
ifeq ($(OS),Windows_NT)
|
||||
$(error Makefile is not supported on windows platform. Please use cmake instead to build sample.)
|
||||
endif
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
|
||||
|
||||
all: vcpy_kernel.code runKernel.hip.out
|
||||
|
||||
runKernel.hip.out: runKernel.cpp
|
||||
$(HIPCC) $(HIPCC_FLAGS) $< -o $@
|
||||
|
||||
vcpy_kernel.code: vcpy_kernel.cpp
|
||||
$(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@
|
||||
|
||||
clean:
|
||||
rm -f *.code *.out
|
||||
@@ -1,161 +0,0 @@
|
||||
/*
|
||||
Copyright (c) 2017 - 2021 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 "hip/hip_runtime.h"
|
||||
#include "hip/hip_runtime_api.h"
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <vector>
|
||||
#include <hip/hip_ext.h>
|
||||
|
||||
#define LEN 64
|
||||
#define SIZE LEN * sizeof(float)
|
||||
|
||||
#define fileName "vcpy_kernel.code"
|
||||
#define HIP_CHECK(cmd) \
|
||||
{ \
|
||||
hipError_t status = cmd; \
|
||||
if (status != hipSuccess) { \
|
||||
std::cout << "error: #" << status << " (" << hipGetErrorString(status) \
|
||||
<< ") at line:" << __LINE__ << ": " << #cmd << std::endl; \
|
||||
abort(); \
|
||||
} \
|
||||
}
|
||||
|
||||
int main() {
|
||||
float *A, *B;
|
||||
float *Ad, *Bd;
|
||||
A = new float[LEN];
|
||||
B = new float[LEN];
|
||||
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
A[i] = i * 1.0f;
|
||||
B[i] = 0.0f;
|
||||
}
|
||||
|
||||
hipInit(0);
|
||||
hipDevice_t device;
|
||||
hipCtx_t context;
|
||||
hipDeviceGet(&device, 0);
|
||||
hipCtxCreate(&context, 0, device);
|
||||
|
||||
hipMalloc((void**)&Ad, SIZE);
|
||||
hipMalloc((void**)&Bd, SIZE);
|
||||
|
||||
hipMemcpyHtoD(hipDeviceptr_t(Ad), A, SIZE);
|
||||
hipMemcpyHtoD((hipDeviceptr_t)(Bd), B, SIZE);
|
||||
hipModule_t Module;
|
||||
HIP_CHECK(hipModuleLoad(&Module, fileName));
|
||||
|
||||
float myDeviceGlobal_h = 42.0;
|
||||
float* deviceGlobal;
|
||||
size_t deviceGlobalSize;
|
||||
HIP_CHECK(hipModuleGetGlobal((void**)&deviceGlobal, &deviceGlobalSize, Module, "myDeviceGlobal"));
|
||||
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(deviceGlobal), &myDeviceGlobal_h, deviceGlobalSize));
|
||||
|
||||
#define ARRAY_SIZE 16
|
||||
|
||||
float myDeviceGlobalArray_h[ARRAY_SIZE];
|
||||
float *myDeviceGlobalArray;
|
||||
size_t myDeviceGlobalArraySize;
|
||||
HIP_CHECK(hipModuleGetGlobal((void**)&myDeviceGlobalArray, &myDeviceGlobalArraySize, Module, "myDeviceGlobalArray"));
|
||||
for (int i = 0; i < ARRAY_SIZE; i++) {
|
||||
myDeviceGlobalArray_h[i] = i * 1000.0f;
|
||||
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(myDeviceGlobalArray), &myDeviceGlobalArray_h, myDeviceGlobalArraySize));
|
||||
}
|
||||
|
||||
struct {
|
||||
void* _Ad;
|
||||
void* _Bd;
|
||||
} args;
|
||||
|
||||
args._Ad = (void*) Ad;
|
||||
args._Bd = (void*) Bd;
|
||||
|
||||
size_t size = sizeof(args);
|
||||
|
||||
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
|
||||
HIP_LAUNCH_PARAM_END};
|
||||
|
||||
{
|
||||
hipFunction_t Function;
|
||||
HIP_CHECK(hipModuleGetFunction(&Function, Module, "hello_world"));
|
||||
HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));
|
||||
|
||||
hipMemcpyDtoH(B, Bd, SIZE);
|
||||
|
||||
int mismatchCount = 0;
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
if (A[i] != B[i]) {
|
||||
mismatchCount++;
|
||||
std::cout << "error: mismatch " << A[i] << " != " << B[i] << std::endl;
|
||||
if (mismatchCount >= 10) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if (mismatchCount == 0) {
|
||||
std::cout << "PASSED!\n";
|
||||
} else {
|
||||
std::cout << "FAILED!\n";
|
||||
};
|
||||
}
|
||||
|
||||
{
|
||||
hipFunction_t Function;
|
||||
HIP_CHECK(hipModuleGetFunction(&Function, Module, "test_globals"));
|
||||
int val =-1;
|
||||
HIP_CHECK(hipFuncGetAttribute(&val, HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES,Function));
|
||||
printf("Shared Size Bytes = %d\n",val);
|
||||
HIP_CHECK(hipFuncGetAttribute(&val, HIP_FUNC_ATTRIBUTE_NUM_REGS, Function));
|
||||
printf("Num Regs = %d\n",val);
|
||||
HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config));
|
||||
|
||||
hipMemcpyDtoH(B, Bd, SIZE);
|
||||
|
||||
int mismatchCount = 0;
|
||||
for (uint32_t i = 0; i < LEN; i++) {
|
||||
float expected = A[i] + myDeviceGlobal_h + myDeviceGlobalArray_h[i % 16];
|
||||
if (expected != B[i]) {
|
||||
mismatchCount++;
|
||||
std::cout << "error: mismatch " << expected << " != " << B[i] << std::endl;
|
||||
if (mismatchCount >= 10) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if (mismatchCount == 0) {
|
||||
std::cout << "PASSED!\n";
|
||||
} else {
|
||||
std::cout << "FAILED!\n";
|
||||
};
|
||||
}
|
||||
|
||||
hipFree(Ad);
|
||||
hipFree(Bd);
|
||||
delete[] A;
|
||||
delete[] B;
|
||||
hipCtxDestroy(context);
|
||||
return 0;
|
||||
}
|
||||
@@ -1,38 +0,0 @@
|
||||
/*
|
||||
Copyright (c) 2017 - 2021 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 "hip/hip_runtime.h"
|
||||
|
||||
#define ARRAY_SIZE (16)
|
||||
|
||||
__device__ float myDeviceGlobal;
|
||||
__device__ float myDeviceGlobalArray[16];
|
||||
|
||||
extern "C" __global__ void hello_world(const float* a, float* b) {
|
||||
int tx = threadIdx.x;
|
||||
b[tx] = a[tx];
|
||||
}
|
||||
|
||||
extern "C" __global__ void test_globals(const float* a, float* b) {
|
||||
int tx = threadIdx.x;
|
||||
b[tx] = a[tx] + myDeviceGlobal + myDeviceGlobalArray[tx % ARRAY_SIZE];
|
||||
}
|
||||
@@ -1,48 +0,0 @@
|
||||
# Copyright (c) 2020 - 2021 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.
|
||||
|
||||
#Follow "README.md" to generate square.cpp if it's missing
|
||||
|
||||
project(square)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# create square.cpp
|
||||
execute_process(COMMAND sh -c "${ROCM_PATH}/hip/bin/hipify-perl ../square.cu > ../square.cpp")
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
# Set compiler and linker
|
||||
set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
|
||||
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})
|
||||
|
||||
# Create the excutable
|
||||
add_executable(square square.cpp)
|
||||
|
||||
# Link with HIP
|
||||
target_link_libraries(square hip::host)
|
||||
@@ -1,47 +0,0 @@
|
||||
# Copyright (c) 2016 - 2021 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.
|
||||
ifeq ($(OS),Windows_NT)
|
||||
$(error Makefile is not supported on windows platform. Please use cmake instead to build sample.)
|
||||
endif
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform)
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
|
||||
ifeq (${HIP_PLATFORM}, nvidia)
|
||||
SOURCES=square.cu
|
||||
else
|
||||
SOURCES=square.cpp
|
||||
endif
|
||||
|
||||
all: square.out
|
||||
|
||||
# Step
|
||||
square.cpp: square.cu
|
||||
$(HIP_PATH)/bin/hipify-perl square.cu > square.cpp
|
||||
|
||||
square.out: $(SOURCES)
|
||||
$(HIPCC) $(CXXFLAGS) $(SOURCES) -o $@
|
||||
|
||||
clean:
|
||||
rm -f *.o *.out square.cpp
|
||||
@@ -1,37 +0,0 @@
|
||||
# Square.md
|
||||
|
||||
Simple test below is an example, shows how to use hipify-perl to port CUDA code to HIP:
|
||||
|
||||
- Add hip/bin path to the PATH
|
||||
|
||||
```
|
||||
$ export PATH=$PATH:[MYHIP]/bin
|
||||
```
|
||||
|
||||
- Define environment variable
|
||||
|
||||
```
|
||||
$ export HIP_PATH=[MYHIP]
|
||||
```
|
||||
|
||||
- Build executible file
|
||||
|
||||
```
|
||||
$ cd ~/hip/samples/0_Intro/square
|
||||
$ make
|
||||
/opt/rocm/hip/bin/hipify-perl square.cu > square.cpp
|
||||
/opt/rocm/hip/bin/hipcc square.cpp -o square.out
|
||||
/opt/rocm/hip/bin/hipcc square.cpp -o square.out
|
||||
```
|
||||
- Execute file
|
||||
```
|
||||
$ ./square.out
|
||||
info: running on device Navi 14 [Radeon Pro W5500]
|
||||
info: allocate host mem ( 7.63 MB)
|
||||
info: allocate device mem ( 7.63 MB)
|
||||
info: copy Host2Device
|
||||
info: launch 'vector_square' kernel
|
||||
info: copy Device2Host
|
||||
info: check result
|
||||
PASSED!
|
||||
```
|
||||
@@ -1,98 +0,0 @@
|
||||
/*
|
||||
Copyright (c) 2015 - 2021 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 <stdio.h>
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#define CHECK(cmd) \
|
||||
{\
|
||||
cudaError_t error = cmd;\
|
||||
if (error != cudaSuccess) { \
|
||||
fprintf(stderr, "error: '%s'(%d) at %s:%d\n", cudaGetErrorString(error), error,__FILE__, __LINE__); \
|
||||
exit(EXIT_FAILURE);\
|
||||
}\
|
||||
}
|
||||
|
||||
|
||||
/*
|
||||
* Square each element in the array A and write to array C.
|
||||
*/
|
||||
template <typename T>
|
||||
__global__ void
|
||||
vector_square(T *C_d, T *A_d, size_t N)
|
||||
{
|
||||
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x ;
|
||||
|
||||
for (size_t i=offset; i<N; i+=stride) {
|
||||
C_d[i] = A_d[i] * A_d[i];
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
int main(int argc, char *argv[])
|
||||
{
|
||||
float *A_d, *C_d;
|
||||
float *A_h, *C_h;
|
||||
size_t N = 1000000;
|
||||
size_t Nbytes = N * sizeof(float);
|
||||
|
||||
cudaDeviceProp props;
|
||||
CHECK(cudaGetDeviceProperties(&props, 0/*deviceID*/));
|
||||
printf ("info: running on device %s\n", props.name);
|
||||
|
||||
printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
|
||||
A_h = (float*)malloc(Nbytes);
|
||||
CHECK(A_h == 0 ? cudaErrorMemoryAllocation : cudaSuccess );
|
||||
C_h = (float*)malloc(Nbytes);
|
||||
CHECK(C_h == 0 ? cudaErrorMemoryAllocation : cudaSuccess );
|
||||
// Fill with Phi + i
|
||||
for (size_t i=0; i<N; i++)
|
||||
{
|
||||
A_h[i] = 1.618f + i;
|
||||
}
|
||||
|
||||
printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0);
|
||||
CHECK(cudaMalloc(&A_d, Nbytes));
|
||||
CHECK(cudaMalloc(&C_d, Nbytes));
|
||||
|
||||
|
||||
printf ("info: copy Host2Device\n");
|
||||
CHECK ( cudaMemcpy(A_d, A_h, Nbytes, cudaMemcpyHostToDevice));
|
||||
|
||||
const unsigned blocks = 512;
|
||||
const unsigned threadsPerBlock = 256;
|
||||
|
||||
printf ("info: launch 'vector_square' kernel\n");
|
||||
vector_square <<<blocks, threadsPerBlock>>> (C_d, A_d, N);
|
||||
|
||||
printf ("info: copy Device2Host\n");
|
||||
CHECK ( cudaMemcpy(C_h, C_d, Nbytes, cudaMemcpyDeviceToHost));
|
||||
|
||||
printf ("info: check result\n");
|
||||
for (size_t i=0; i<N; i++) {
|
||||
if (C_h[i] != A_h[i] * A_h[i]) {
|
||||
CHECK(cudaErrorUnknown);
|
||||
}
|
||||
}
|
||||
printf ("PASSED!\n");
|
||||
}
|
||||
@@ -1,96 +0,0 @@
|
||||
/*
|
||||
Copyright (c) 2015 - 2021 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 <stdio.h>
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
#define CHECK(cmd) \
|
||||
{ \
|
||||
hipError_t error = cmd; \
|
||||
if (error != hipSuccess) { \
|
||||
fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error, \
|
||||
__FILE__, __LINE__); \
|
||||
exit(EXIT_FAILURE); \
|
||||
} \
|
||||
}
|
||||
|
||||
/*
|
||||
* Square each element in the array A and write to array C.
|
||||
*/
|
||||
template <typename T>
|
||||
__global__ void vector_square(T* C_d, const T* A_d, size_t N) {
|
||||
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
|
||||
for (size_t i = offset; i < N; i += stride) {
|
||||
C_d[i] = A_d[i] * A_d[i];
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
int main(int argc, char* argv[]) {
|
||||
float *A_d, *C_d;
|
||||
float *A_h, *C_h;
|
||||
size_t N = 1000000;
|
||||
size_t Nbytes = N * sizeof(float);
|
||||
static int device = 0;
|
||||
CHECK(hipSetDevice(device));
|
||||
hipDeviceProp_t props;
|
||||
CHECK(hipGetDeviceProperties(&props, device /*deviceID*/));
|
||||
printf("info: running on device %s\n", props.name);
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
printf("info: architecture on AMD GPU device is: %s\n", props.gcnArchName);
|
||||
#endif
|
||||
printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0);
|
||||
A_h = (float*)malloc(Nbytes);
|
||||
CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
C_h = (float*)malloc(Nbytes);
|
||||
CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
// Fill with Phi + i
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
A_h[i] = 1.618f + i;
|
||||
}
|
||||
|
||||
printf("info: allocate device mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0);
|
||||
CHECK(hipMalloc(&A_d, Nbytes));
|
||||
CHECK(hipMalloc(&C_d, Nbytes));
|
||||
|
||||
printf("info: copy Host2Device\n");
|
||||
CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
const unsigned blocks = 512;
|
||||
const unsigned threadsPerBlock = 256;
|
||||
|
||||
printf("info: launch 'vector_square' kernel\n");
|
||||
hipLaunchKernelGGL(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
|
||||
|
||||
printf("info: copy Device2Host\n");
|
||||
CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
|
||||
printf("info: check result\n");
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
if (C_h[i] != A_h[i] * A_h[i]) {
|
||||
CHECK(hipErrorUnknown);
|
||||
}
|
||||
}
|
||||
printf("PASSED!\n");
|
||||
}
|
||||
@@ -1,44 +0,0 @@
|
||||
# Copyright (c) 2020 - 2021 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.
|
||||
|
||||
project(hipBusBandwidth)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
# Set compiler and linker
|
||||
set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
|
||||
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})
|
||||
set(CMAKE_BUILD_TYPE Release)
|
||||
|
||||
# Create the excutable
|
||||
add_executable(hipBusBandwidth hipBusBandwidth.cpp ResultDatabase.cpp)
|
||||
|
||||
# Link with HIP
|
||||
target_link_libraries(hipBusBandwidth hip::host)
|
||||
@@ -1,27 +0,0 @@
|
||||
|
||||
Copyright (c) 2011, UT-Battelle, LLC
|
||||
All rights reserved.
|
||||
|
||||
Redistribution and use in source and binary forms, with or without
|
||||
modification, are permitted provided that the following conditions are met:
|
||||
|
||||
* Redistributions of source code must retain the above copyright
|
||||
notice, this list of conditions and the following disclaimer.
|
||||
* Redistributions in binary form must reproduce the above copyright
|
||||
notice, this list of conditions and the following disclaimer in the
|
||||
documentation and/or other materials provided with the distribution.
|
||||
* Neither the name of Oak Ridge National Laboratory, nor UT-Battelle, LLC, nor
|
||||
the names of its contributors may be used to endorse or promote products
|
||||
derived from this software without specific prior written permission.
|
||||
|
||||
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
|
||||
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
|
||||
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
|
||||
DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
|
||||
FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
|
||||
DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
|
||||
SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
|
||||
CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
|
||||
OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
|
||||
OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
|
||||
@@ -1,43 +0,0 @@
|
||||
# Copyright (c) 2016 - 2021 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.
|
||||
ifeq ($(OS),Windows_NT)
|
||||
$(error Makefile is not supported on windows platform. Please use cmake instead to build sample.)
|
||||
endif
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
|
||||
EXE=hipBusBandwidth
|
||||
CXXFLAGS = -O3
|
||||
|
||||
all: install
|
||||
|
||||
$(EXE): hipBusBandwidth.cpp ResultDatabase.cpp
|
||||
$(HIPCC) $(CXXFLAGS) $^ -o $@
|
||||
|
||||
install: $(EXE)
|
||||
cp $(EXE) $(HIP_PATH)/bin
|
||||
|
||||
|
||||
clean:
|
||||
rm -f *.o $(EXE)
|
||||
@@ -1,462 +0,0 @@
|
||||
#include "ResultDatabase.h"
|
||||
|
||||
#include <cfloat>
|
||||
#include <algorithm>
|
||||
#include <cmath>
|
||||
#include <iomanip>
|
||||
|
||||
using namespace std;
|
||||
|
||||
#define SORT_RETAIN_ATTS_ORDER 1
|
||||
|
||||
|
||||
bool ResultDatabase::Result::operator<(const Result& rhs) const {
|
||||
if (test < rhs.test) return true;
|
||||
if (test > rhs.test) return false;
|
||||
#if (SORT_RETAIN_ATTS_ORDER == 0)
|
||||
// For ties, sort by the value of the attribute:
|
||||
if (atts < rhs.atts) return true;
|
||||
if (atts > rhs.atts) return false;
|
||||
#endif
|
||||
return false; // less-operator returns false on equal
|
||||
}
|
||||
|
||||
double ResultDatabase::Result::GetMin() const {
|
||||
double r = FLT_MAX;
|
||||
for (int i = 0; i < value.size(); i++) {
|
||||
r = min(r, value[i]);
|
||||
}
|
||||
return r;
|
||||
}
|
||||
|
||||
double ResultDatabase::Result::GetMax() const {
|
||||
double r = -FLT_MAX;
|
||||
for (int i = 0; i < value.size(); i++) {
|
||||
r = max(r, value[i]);
|
||||
}
|
||||
return r;
|
||||
}
|
||||
|
||||
double ResultDatabase::Result::GetMedian() const { return GetPercentile(50); }
|
||||
|
||||
double ResultDatabase::Result::GetPercentile(double q) const {
|
||||
int n = value.size();
|
||||
if (n == 0) return FLT_MAX;
|
||||
if (n == 1) return value[0];
|
||||
|
||||
if (q <= 0) return value[0];
|
||||
if (q >= 100) return value[n - 1];
|
||||
|
||||
double index = ((n + 1.) * q / 100.) - 1;
|
||||
|
||||
vector<double> sorted = value;
|
||||
sort(sorted.begin(), sorted.end());
|
||||
|
||||
if (n == 2) return (sorted[0] * (1 - q / 100.) + sorted[1] * (q / 100.));
|
||||
|
||||
int index_lo = int(index);
|
||||
double frac = index - index_lo;
|
||||
if (frac == 0) return sorted[index_lo];
|
||||
|
||||
double lo = sorted[index_lo];
|
||||
double hi = sorted[index_lo + 1];
|
||||
return lo + (hi - lo) * frac;
|
||||
}
|
||||
|
||||
double ResultDatabase::Result::GetMean() const {
|
||||
double r = 0;
|
||||
for (int i = 0; i < value.size(); i++) {
|
||||
r += value[i];
|
||||
}
|
||||
return r / double(value.size());
|
||||
}
|
||||
|
||||
double ResultDatabase::Result::GetStdDev() const {
|
||||
double r = 0;
|
||||
double u = GetMean();
|
||||
if (u == FLT_MAX) return FLT_MAX;
|
||||
for (int i = 0; i < value.size(); i++) {
|
||||
r += (value[i] - u) * (value[i] - u);
|
||||
}
|
||||
r = sqrt(r / value.size());
|
||||
return r;
|
||||
}
|
||||
|
||||
|
||||
void ResultDatabase::AddResults(const string& test, const string& atts, const string& unit,
|
||||
const vector<double>& values) {
|
||||
for (int i = 0; i < values.size(); i++) {
|
||||
AddResult(test, atts, unit, values[i]);
|
||||
}
|
||||
}
|
||||
|
||||
static string RemoveAllButLeadingSpaces(const string& a) {
|
||||
string b;
|
||||
int n = a.length();
|
||||
int i = 0;
|
||||
while (i < n && a[i] == ' ') {
|
||||
b += a[i];
|
||||
++i;
|
||||
}
|
||||
for (; i < n; i++) {
|
||||
if (a[i] != ' ' && a[i] != '\t') b += a[i];
|
||||
}
|
||||
return b;
|
||||
}
|
||||
|
||||
void ResultDatabase::AddResult(const string& test_orig, const string& atts_orig,
|
||||
const string& unit_orig, double value) {
|
||||
string test = RemoveAllButLeadingSpaces(test_orig);
|
||||
string atts = RemoveAllButLeadingSpaces(atts_orig);
|
||||
string unit = RemoveAllButLeadingSpaces(unit_orig);
|
||||
int index;
|
||||
for (index = 0; index < results.size(); index++) {
|
||||
if (results[index].test == test && results[index].atts == atts) {
|
||||
if (results[index].unit != unit) throw "Internal error: mixed units";
|
||||
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
if (index >= results.size()) {
|
||||
Result r;
|
||||
r.test = test;
|
||||
r.atts = atts;
|
||||
r.unit = unit;
|
||||
results.push_back(r);
|
||||
}
|
||||
|
||||
results[index].value.push_back(value);
|
||||
}
|
||||
|
||||
// ****************************************************************************
|
||||
// Method: ResultDatabase::DumpDetailed
|
||||
//
|
||||
// Purpose:
|
||||
// Writes the full results, including all trials.
|
||||
//
|
||||
// Arguments:
|
||||
// out where to print
|
||||
//
|
||||
// Programmer: Jeremy Meredith
|
||||
// Creation: August 14, 2009
|
||||
//
|
||||
// Modifications:
|
||||
// Jeremy Meredith, Wed Nov 10 14:25:17 EST 2010
|
||||
// Renamed to DumpDetailed to make room for a DumpSummary.
|
||||
//
|
||||
// Jeremy Meredith, Thu Nov 11 11:39:57 EST 2010
|
||||
// Added note about (*) missing value tag.
|
||||
//
|
||||
// Jeremy Meredith, Tue Nov 23 13:57:02 EST 2010
|
||||
// Changed note about missing values to be worded a little better.
|
||||
//
|
||||
// ****************************************************************************
|
||||
void ResultDatabase::DumpDetailed(ostream& out) {
|
||||
vector<Result> sorted(results);
|
||||
|
||||
stable_sort(sorted.begin(), sorted.end());
|
||||
|
||||
const int testNameW = 24;
|
||||
const int attW = 12;
|
||||
const int fieldW = 11;
|
||||
out << std::fixed << right << std::setprecision(4);
|
||||
|
||||
int maxtrials = 1;
|
||||
for (int i = 0; i < sorted.size(); i++) {
|
||||
if (sorted[i].value.size() > maxtrials) maxtrials = sorted[i].value.size();
|
||||
}
|
||||
|
||||
// TODO: in big parallel runs, the "trials" are the procs
|
||||
// and we really don't want to print them all out....
|
||||
out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "median\t"
|
||||
<< "mean\t"
|
||||
<< "stddev\t"
|
||||
<< "min\t"
|
||||
<< "max\t";
|
||||
for (int i = 0; i < maxtrials; i++) out << "trial" << i << "\t";
|
||||
out << endl;
|
||||
|
||||
for (int i = 0; i < sorted.size(); i++) {
|
||||
Result& r = sorted[i];
|
||||
out << setw(testNameW) << r.test + "\t";
|
||||
out << setw(attW) << r.atts + "\t";
|
||||
out << setw(fieldW) << r.unit + "\t";
|
||||
if (r.GetMedian() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMedian() << "\t";
|
||||
if (r.GetMean() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMean() << "\t";
|
||||
if (r.GetStdDev() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetStdDev() << "\t";
|
||||
if (r.GetMin() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMin() << "\t";
|
||||
if (r.GetMax() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMax() << "\t";
|
||||
for (int j = 0; j < r.value.size(); j++) {
|
||||
if (r.value[j] == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.value[j] << "\t";
|
||||
}
|
||||
|
||||
out << endl;
|
||||
}
|
||||
out << endl
|
||||
<< "Note: Any results marked with (*) had missing values." << endl
|
||||
<< " This can occur on systems with a mixture of" << endl
|
||||
<< " device types or architectural capabilities." << endl;
|
||||
}
|
||||
|
||||
|
||||
// ****************************************************************************
|
||||
// Method: ResultDatabase::DumpDetailed
|
||||
//
|
||||
// Purpose:
|
||||
// Writes the summary results (min/max/stddev/med/mean), but not
|
||||
// every individual trial.
|
||||
//
|
||||
// Arguments:
|
||||
// out where to print
|
||||
//
|
||||
// Programmer: Jeremy Meredith
|
||||
// Creation: November 10, 2010
|
||||
//
|
||||
// Modifications:
|
||||
// Jeremy Meredith, Thu Nov 11 11:39:57 EST 2010
|
||||
// Added note about (*) missing value tag.
|
||||
//
|
||||
// ****************************************************************************
|
||||
void ResultDatabase::DumpSummary(ostream& out) {
|
||||
vector<Result> sorted(results);
|
||||
|
||||
stable_sort(sorted.begin(), sorted.end());
|
||||
|
||||
const int testNameW = 24;
|
||||
const int attW = 12;
|
||||
const int fieldW = 9;
|
||||
out << std::fixed << right << std::setprecision(4);
|
||||
|
||||
// TODO: in big parallel runs, the "trials" are the procs
|
||||
// and we really don't want to print them all out....
|
||||
out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "units\t"
|
||||
<< "median\t"
|
||||
<< "mean\t"
|
||||
<< "stddev\t"
|
||||
<< "min\t"
|
||||
<< "max\t";
|
||||
out << endl;
|
||||
|
||||
for (int i = 0; i < sorted.size(); i++) {
|
||||
Result& r = sorted[i];
|
||||
out << setw(testNameW) << r.test + "\t";
|
||||
out << setw(attW) << r.atts + "\t";
|
||||
out << setw(fieldW) << r.unit + "\t";
|
||||
if (r.GetMedian() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMedian() << "\t";
|
||||
if (r.GetMean() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMean() << "\t";
|
||||
if (r.GetStdDev() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetStdDev() << "\t";
|
||||
if (r.GetMin() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMin() << "\t";
|
||||
if (r.GetMax() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMax() << "\t";
|
||||
|
||||
out << endl;
|
||||
}
|
||||
out << endl
|
||||
<< "Note: results marked with (*) had missing values such as" << endl
|
||||
<< "might occur with a mixture of architectural capabilities." << endl;
|
||||
}
|
||||
|
||||
// ****************************************************************************
|
||||
// Method: ResultDatabase::ClearAllResults
|
||||
//
|
||||
// Purpose:
|
||||
// Clears all existing results from the ResultDatabase; used for multiple passes
|
||||
// of the same test or multiple tests.
|
||||
//
|
||||
// Arguments:
|
||||
//
|
||||
// Programmer: Jeffrey Young
|
||||
// Creation: September 10th, 2014
|
||||
//
|
||||
// Modifications:
|
||||
//
|
||||
//
|
||||
// ****************************************************************************
|
||||
void ResultDatabase::ClearAllResults() { results.clear(); }
|
||||
|
||||
// ****************************************************************************
|
||||
// Method: ResultDatabase::DumpCsv
|
||||
//
|
||||
// Purpose:
|
||||
// Writes either detailed or summary results (min/max/stddev/med/mean), but not
|
||||
// every individual trial.
|
||||
//
|
||||
// Arguments:
|
||||
// out file to print CSV results
|
||||
//
|
||||
// Programmer: Jeffrey Young
|
||||
// Creation: August 28th, 2014
|
||||
//
|
||||
// Modifications:
|
||||
//
|
||||
// ****************************************************************************
|
||||
void ResultDatabase::DumpCsv(string fileName) {
|
||||
bool emptyFile;
|
||||
vector<Result> sorted(results);
|
||||
|
||||
stable_sort(sorted.begin(), sorted.end());
|
||||
|
||||
// Check to see if the file is empty - if so, add the headers
|
||||
emptyFile = this->IsFileEmpty(fileName);
|
||||
|
||||
// Open file and append by default
|
||||
ofstream out;
|
||||
out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app);
|
||||
|
||||
// Add headers only for empty files
|
||||
if (emptyFile) {
|
||||
// TODO: in big parallel runs, the "trials" are the procs
|
||||
// and we really don't want to print them all out....
|
||||
out << "test, "
|
||||
<< "atts, "
|
||||
<< "units, "
|
||||
<< "median, "
|
||||
<< "mean, "
|
||||
<< "stddev, "
|
||||
<< "min, "
|
||||
<< "max, ";
|
||||
out << endl;
|
||||
}
|
||||
|
||||
for (int i = 0; i < sorted.size(); i++) {
|
||||
Result& r = sorted[i];
|
||||
out << r.test << ", ";
|
||||
out << r.atts << ", ";
|
||||
out << r.unit << ", ";
|
||||
if (r.GetMedian() == FLT_MAX)
|
||||
out << "N/A, ";
|
||||
else
|
||||
out << r.GetMedian() << ", ";
|
||||
if (r.GetMean() == FLT_MAX)
|
||||
out << "N/A, ";
|
||||
else
|
||||
out << r.GetMean() << ", ";
|
||||
if (r.GetStdDev() == FLT_MAX)
|
||||
out << "N/A, ";
|
||||
else
|
||||
out << r.GetStdDev() << ", ";
|
||||
if (r.GetMin() == FLT_MAX)
|
||||
out << "N/A, ";
|
||||
else
|
||||
out << r.GetMin() << ", ";
|
||||
if (r.GetMax() == FLT_MAX)
|
||||
out << "N/A, ";
|
||||
else
|
||||
out << r.GetMax() << ", ";
|
||||
|
||||
out << endl;
|
||||
}
|
||||
out << endl;
|
||||
|
||||
out.close();
|
||||
}
|
||||
|
||||
// ****************************************************************************
|
||||
// Method: ResultDatabase::IsFileEmpty
|
||||
//
|
||||
// Purpose:
|
||||
// Returns whether a file is empty - used as a helper for CSV printing
|
||||
//
|
||||
// Arguments:
|
||||
// file The input file to check for emptiness
|
||||
//
|
||||
// Programmer: Jeffrey Young
|
||||
// Creation: August 28th, 2014
|
||||
//
|
||||
// Modifications:
|
||||
//
|
||||
// ****************************************************************************
|
||||
|
||||
bool ResultDatabase::IsFileEmpty(string fileName) {
|
||||
|
||||
ifstream file(fileName.c_str());
|
||||
|
||||
// If the file doesn't exist it is by definition empty
|
||||
if (!file.good()) {
|
||||
return true;
|
||||
} else {
|
||||
bool fileEmpty;
|
||||
fileEmpty = (bool)(file.peek() == ifstream::traits_type::eof());
|
||||
file.close();
|
||||
|
||||
return fileEmpty;
|
||||
}
|
||||
|
||||
// Otherwise, return false
|
||||
return false;
|
||||
}
|
||||
|
||||
|
||||
// ****************************************************************************
|
||||
// Method: ResultDatabase::GetResultsForTest
|
||||
//
|
||||
// Purpose:
|
||||
// Returns a vector of results for just one test name.
|
||||
//
|
||||
// Arguments:
|
||||
// test the name of the test results to search for
|
||||
//
|
||||
// Programmer: Jeremy Meredith
|
||||
// Creation: December 3, 2010
|
||||
//
|
||||
// Modifications:
|
||||
//
|
||||
// ****************************************************************************
|
||||
vector<ResultDatabase::Result> ResultDatabase::GetResultsForTest(const string& test) {
|
||||
// get only the given test results
|
||||
vector<Result> retval;
|
||||
for (int i = 0; i < results.size(); i++) {
|
||||
Result& r = results[i];
|
||||
if (r.test == test) retval.push_back(r);
|
||||
}
|
||||
return retval;
|
||||
}
|
||||
|
||||
// ****************************************************************************
|
||||
// Method: ResultDatabase::GetResults
|
||||
//
|
||||
// Purpose:
|
||||
// Returns all the results.
|
||||
//
|
||||
// Arguments:
|
||||
//
|
||||
// Programmer: Jeremy Meredith
|
||||
// Creation: December 3, 2010
|
||||
//
|
||||
// Modifications:
|
||||
//
|
||||
// ****************************************************************************
|
||||
const vector<ResultDatabase::Result>& ResultDatabase::GetResults() const { return results; }
|
||||
@@ -1,89 +0,0 @@
|
||||
#ifndef RESULT_DATABASE_H
|
||||
#define RESULT_DATABASE_H
|
||||
|
||||
#include <string>
|
||||
#include <vector>
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <cfloat>
|
||||
using std::ifstream;
|
||||
using std::ofstream;
|
||||
using std::ostream;
|
||||
using std::string;
|
||||
using std::vector;
|
||||
|
||||
|
||||
// ****************************************************************************
|
||||
// Class: ResultDatabase
|
||||
//
|
||||
// Purpose:
|
||||
// Track numerical results as they are generated.
|
||||
// Print statistics of raw results.
|
||||
//
|
||||
// Programmer: Jeremy Meredith
|
||||
// Creation: June 12, 2009
|
||||
//
|
||||
// Modifications:
|
||||
// Jeremy Meredith, Wed Nov 10 14:20:47 EST 2010
|
||||
// Split timing reports into detailed and summary. E.g. for serial code,
|
||||
// we might report all trial values, but skip them in parallel.
|
||||
//
|
||||
// Jeremy Meredith, Thu Nov 11 11:40:18 EST 2010
|
||||
// Added check for missing value tag.
|
||||
//
|
||||
// Jeremy Meredith, Mon Nov 22 13:37:10 EST 2010
|
||||
// Added percentile statistic.
|
||||
//
|
||||
// Jeremy Meredith, Fri Dec 3 16:30:31 EST 2010
|
||||
// Added a method to extract a subset of results based on test name. Also,
|
||||
// the Result class is now public, so that clients can use them directly.
|
||||
// Added a GetResults method as well, and made several functions const.
|
||||
//
|
||||
// ****************************************************************************
|
||||
class ResultDatabase {
|
||||
public:
|
||||
//
|
||||
// A performance result for a single SHOC benchmark run.
|
||||
//
|
||||
struct Result {
|
||||
string test; // e.g. "readback"
|
||||
string atts; // e.g. "pagelocked 4k^2"
|
||||
string unit; // e.g. "MB/sec"
|
||||
vector<double> value; // e.g. "837.14"
|
||||
double GetMin() const;
|
||||
double GetMax() const;
|
||||
double GetMedian() const;
|
||||
double GetPercentile(double q) const;
|
||||
double GetMean() const;
|
||||
double GetStdDev() const;
|
||||
|
||||
bool operator<(const Result& rhs) const;
|
||||
|
||||
bool HadAnyFLTMAXValues() const {
|
||||
for (int i = 0; i < value.size(); ++i) {
|
||||
if (value[i] >= FLT_MAX) return true;
|
||||
}
|
||||
return false;
|
||||
}
|
||||
};
|
||||
|
||||
protected:
|
||||
vector<Result> results;
|
||||
|
||||
public:
|
||||
void AddResult(const string& test, const string& atts, const string& unit, double value);
|
||||
void AddResults(const string& test, const string& atts, const string& unit,
|
||||
const vector<double>& values);
|
||||
vector<Result> GetResultsForTest(const string& test);
|
||||
const vector<Result>& GetResults() const;
|
||||
void ClearAllResults();
|
||||
void DumpDetailed(ostream&);
|
||||
void DumpSummary(ostream&);
|
||||
void DumpCsv(string fileName);
|
||||
|
||||
private:
|
||||
bool IsFileEmpty(string fileName);
|
||||
};
|
||||
|
||||
|
||||
#endif
|
||||
Rozdílový obsah nebyl zobrazen, protože je příliš veliký
Načíst rozdílové porovnání
@@ -1,55 +0,0 @@
|
||||
# Copyright (c) 2020 - 2021 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.
|
||||
|
||||
project(hipCommander)
|
||||
|
||||
cmake_minimum_required(VERSION 3.10)
|
||||
|
||||
if (NOT DEFINED ROCM_PATH )
|
||||
set ( ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory." )
|
||||
endif ()
|
||||
|
||||
# Search for rocm in common locations
|
||||
list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH})
|
||||
|
||||
# Find hip
|
||||
find_package(hip)
|
||||
|
||||
# Set compiler and linker
|
||||
set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
|
||||
set(CMAKE_CXX_LINKER ${HIP_HIPCC_EXECUTABLE})
|
||||
set(CMAKE_BUILD_TYPE Release)
|
||||
|
||||
# Create the excutable
|
||||
add_executable(hipCommander hipCommander.cpp)
|
||||
|
||||
# Generate code object
|
||||
add_custom_target(
|
||||
codeobj
|
||||
ALL
|
||||
COMMAND ${HIP_HIPCC_EXECUTABLE} --genco ../nullkernel.hip.cpp -o nullkernel.hsaco
|
||||
COMMENT "codeobj generated"
|
||||
)
|
||||
|
||||
add_dependencies(hipCommander codeobj)
|
||||
|
||||
# Link with HIP
|
||||
target_link_libraries(hipCommander hip::host)
|
||||
set_property(TARGET hipCommander PROPERTY CXX_STANDARD 11)
|
||||
@@ -1,27 +0,0 @@
|
||||
|
||||
Copyright (c) 2011, UT-Battelle, LLC
|
||||
All rights reserved.
|
||||
|
||||
Redistribution and use in source and binary forms, with or without
|
||||
modification, are permitted provided that the following conditions are met:
|
||||
|
||||
* Redistributions of source code must retain the above copyright
|
||||
notice, this list of conditions and the following disclaimer.
|
||||
* Redistributions in binary form must reproduce the above copyright
|
||||
notice, this list of conditions and the following disclaimer in the
|
||||
documentation and/or other materials provided with the distribution.
|
||||
* Neither the name of Oak Ridge National Laboratory, nor UT-Battelle, LLC, nor
|
||||
the names of its contributors may be used to endorse or promote products
|
||||
derived from this software without specific prior written permission.
|
||||
|
||||
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
|
||||
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
|
||||
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
|
||||
DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
|
||||
FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
|
||||
DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
|
||||
SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
|
||||
CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
|
||||
OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
|
||||
OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
|
||||
@@ -1,53 +0,0 @@
|
||||
# Copyright (c) 2016 - 2021 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.
|
||||
ifeq ($(OS),Windows_NT)
|
||||
$(error Makefile is not supported on windows platform. Please use cmake instead to build sample.)
|
||||
endif
|
||||
ROCM_PATH?= $(wildcard /opt/rocm/)
|
||||
HIP_PATH?= $(wildcard $(ROCM_PATH)/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
|
||||
EXE=hipCommander
|
||||
OPT=-O3
|
||||
#CXXFLAGS = -O3 -g
|
||||
CXXFLAGS = $(OPT) --std=c++11
|
||||
|
||||
HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform)
|
||||
|
||||
CODE_OBJECTS=nullkernel.hsaco
|
||||
|
||||
all: ${EXE} ${CODE_OBJECTS}
|
||||
|
||||
$(EXE): hipCommander.cpp
|
||||
$(HIPCC) $(CXXFLAGS) $^ -o $@
|
||||
|
||||
nullkernel.hsaco : nullkernel.hip.cpp
|
||||
$(HIPCC) --genco nullkernel.hip.cpp -o nullkernel.hsaco
|
||||
|
||||
|
||||
install: $(EXE)
|
||||
cp $(EXE) $(HIP_PATH)/bin
|
||||
|
||||
|
||||
clean:
|
||||
rm -f *.o *.co $(EXE)
|
||||
@@ -1,454 +0,0 @@
|
||||
#include "ResultDatabase.h"
|
||||
|
||||
#include <cfloat>
|
||||
#include <algorithm>
|
||||
#include <cmath>
|
||||
#include <iomanip>
|
||||
|
||||
using namespace std;
|
||||
|
||||
bool ResultDatabase::Result::operator<(const Result& rhs) const {
|
||||
if (test < rhs.test) return true;
|
||||
if (test > rhs.test) return false;
|
||||
if (atts < rhs.atts) return true;
|
||||
if (atts > rhs.atts) return false;
|
||||
return false; // less-operator returns false on equal
|
||||
}
|
||||
|
||||
double ResultDatabase::Result::GetMin() const {
|
||||
double r = FLT_MAX;
|
||||
for (int i = 0; i < value.size(); i++) {
|
||||
r = min(r, value[i]);
|
||||
}
|
||||
return r;
|
||||
}
|
||||
|
||||
double ResultDatabase::Result::GetMax() const {
|
||||
double r = -FLT_MAX;
|
||||
for (int i = 0; i < value.size(); i++) {
|
||||
r = max(r, value[i]);
|
||||
}
|
||||
return r;
|
||||
}
|
||||
|
||||
double ResultDatabase::Result::GetMedian() const { return GetPercentile(50); }
|
||||
|
||||
double ResultDatabase::Result::GetPercentile(double q) const {
|
||||
int n = value.size();
|
||||
if (n == 0) return FLT_MAX;
|
||||
if (n == 1) return value[0];
|
||||
|
||||
if (q <= 0) return value[0];
|
||||
if (q >= 100) return value[n - 1];
|
||||
|
||||
double index = ((n + 1.) * q / 100.) - 1;
|
||||
|
||||
vector<double> sorted = value;
|
||||
sort(sorted.begin(), sorted.end());
|
||||
|
||||
if (n == 2) return (sorted[0] * (1 - q / 100.) + sorted[1] * (q / 100.));
|
||||
|
||||
int index_lo = int(index);
|
||||
double frac = index - index_lo;
|
||||
if (frac == 0) return sorted[index_lo];
|
||||
|
||||
double lo = sorted[index_lo];
|
||||
double hi = sorted[index_lo + 1];
|
||||
return lo + (hi - lo) * frac;
|
||||
}
|
||||
|
||||
double ResultDatabase::Result::GetMean() const {
|
||||
double r = 0;
|
||||
for (int i = 0; i < value.size(); i++) {
|
||||
r += value[i];
|
||||
}
|
||||
return r / double(value.size());
|
||||
}
|
||||
|
||||
double ResultDatabase::Result::GetStdDev() const {
|
||||
double r = 0;
|
||||
double u = GetMean();
|
||||
if (u == FLT_MAX) return FLT_MAX;
|
||||
for (int i = 0; i < value.size(); i++) {
|
||||
r += (value[i] - u) * (value[i] - u);
|
||||
}
|
||||
r = sqrt(r / value.size());
|
||||
return r;
|
||||
}
|
||||
|
||||
|
||||
void ResultDatabase::AddResults(const string& test, const string& atts, const string& unit,
|
||||
const vector<double>& values) {
|
||||
for (int i = 0; i < values.size(); i++) {
|
||||
AddResult(test, atts, unit, values[i]);
|
||||
}
|
||||
}
|
||||
|
||||
static string RemoveAllButLeadingSpaces(const string& a) {
|
||||
string b;
|
||||
int n = a.length();
|
||||
int i = 0;
|
||||
while (i < n && a[i] == ' ') {
|
||||
b += a[i];
|
||||
++i;
|
||||
}
|
||||
for (; i < n; i++) {
|
||||
if (a[i] != ' ' && a[i] != '\t') b += a[i];
|
||||
}
|
||||
return b;
|
||||
}
|
||||
|
||||
void ResultDatabase::AddResult(const string& test_orig, const string& atts_orig,
|
||||
const string& unit_orig, double value) {
|
||||
string test = RemoveAllButLeadingSpaces(test_orig);
|
||||
string atts = RemoveAllButLeadingSpaces(atts_orig);
|
||||
string unit = RemoveAllButLeadingSpaces(unit_orig);
|
||||
int index;
|
||||
for (index = 0; index < results.size(); index++) {
|
||||
if (results[index].test == test && results[index].atts == atts) {
|
||||
if (results[index].unit != unit) throw "Internal error: mixed units";
|
||||
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
if (index >= results.size()) {
|
||||
Result r;
|
||||
r.test = test;
|
||||
r.atts = atts;
|
||||
r.unit = unit;
|
||||
results.push_back(r);
|
||||
}
|
||||
|
||||
results[index].value.push_back(value);
|
||||
}
|
||||
|
||||
// ****************************************************************************
|
||||
// Method: ResultDatabase::DumpDetailed
|
||||
//
|
||||
// Purpose:
|
||||
// Writes the full results, including all trials.
|
||||
//
|
||||
// Arguments:
|
||||
// out where to print
|
||||
//
|
||||
// Programmer: Jeremy Meredith
|
||||
// Creation: August 14, 2009
|
||||
//
|
||||
// Modifications:
|
||||
// Jeremy Meredith, Wed Nov 10 14:25:17 EST 2010
|
||||
// Renamed to DumpDetailed to make room for a DumpSummary.
|
||||
//
|
||||
// Jeremy Meredith, Thu Nov 11 11:39:57 EST 2010
|
||||
// Added note about (*) missing value tag.
|
||||
//
|
||||
// Jeremy Meredith, Tue Nov 23 13:57:02 EST 2010
|
||||
// Changed note about missing values to be worded a little better.
|
||||
//
|
||||
// ****************************************************************************
|
||||
void ResultDatabase::DumpDetailed(ostream& out) {
|
||||
vector<Result> sorted(results);
|
||||
sort(sorted.begin(), sorted.end());
|
||||
|
||||
const int testNameW = 24;
|
||||
const int attW = 12;
|
||||
const int fieldW = 11;
|
||||
out << std::fixed << right << std::setprecision(4);
|
||||
|
||||
int maxtrials = 1;
|
||||
for (int i = 0; i < sorted.size(); i++) {
|
||||
if (sorted[i].value.size() > maxtrials) maxtrials = sorted[i].value.size();
|
||||
}
|
||||
|
||||
// TODO: in big parallel runs, the "trials" are the procs
|
||||
// and we really don't want to print them all out....
|
||||
out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "median\t"
|
||||
<< "mean\t"
|
||||
<< "stddev\t"
|
||||
<< "min\t"
|
||||
<< "max\t";
|
||||
for (int i = 0; i < maxtrials; i++) out << "trial" << i << "\t";
|
||||
out << endl;
|
||||
|
||||
for (int i = 0; i < sorted.size(); i++) {
|
||||
Result& r = sorted[i];
|
||||
out << setw(testNameW) << r.test + "\t";
|
||||
out << setw(attW) << r.atts + "\t";
|
||||
out << setw(fieldW) << r.unit + "\t";
|
||||
if (r.GetMedian() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMedian() << "\t";
|
||||
if (r.GetMean() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMean() << "\t";
|
||||
if (r.GetStdDev() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetStdDev() << "\t";
|
||||
if (r.GetMin() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMin() << "\t";
|
||||
if (r.GetMax() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMax() << "\t";
|
||||
for (int j = 0; j < r.value.size(); j++) {
|
||||
if (r.value[j] == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.value[j] << "\t";
|
||||
}
|
||||
|
||||
out << endl;
|
||||
}
|
||||
out << endl
|
||||
<< "Note: Any results marked with (*) had missing values." << endl
|
||||
<< " This can occur on systems with a mixture of" << endl
|
||||
<< " device types or architectural capabilities." << endl;
|
||||
}
|
||||
|
||||
|
||||
// ****************************************************************************
|
||||
// Method: ResultDatabase::DumpDetailed
|
||||
//
|
||||
// Purpose:
|
||||
// Writes the summary results (min/max/stddev/med/mean), but not
|
||||
// every individual trial.
|
||||
//
|
||||
// Arguments:
|
||||
// out where to print
|
||||
//
|
||||
// Programmer: Jeremy Meredith
|
||||
// Creation: November 10, 2010
|
||||
//
|
||||
// Modifications:
|
||||
// Jeremy Meredith, Thu Nov 11 11:39:57 EST 2010
|
||||
// Added note about (*) missing value tag.
|
||||
//
|
||||
// ****************************************************************************
|
||||
void ResultDatabase::DumpSummary(ostream& out) {
|
||||
vector<Result> sorted(results);
|
||||
sort(sorted.begin(), sorted.end());
|
||||
|
||||
const int testNameW = 24;
|
||||
const int attW = 12;
|
||||
const int fieldW = 9;
|
||||
out << std::fixed << right << std::setprecision(4);
|
||||
|
||||
// TODO: in big parallel runs, the "trials" are the procs
|
||||
// and we really don't want to print them all out....
|
||||
out << setw(testNameW) << "test\t" << setw(attW) << "atts\t" << setw(fieldW) << "units\t"
|
||||
<< "median\t"
|
||||
<< "mean\t"
|
||||
<< "stddev\t"
|
||||
<< "min\t"
|
||||
<< "max\t";
|
||||
out << endl;
|
||||
|
||||
for (int i = 0; i < sorted.size(); i++) {
|
||||
Result& r = sorted[i];
|
||||
out << setw(testNameW) << r.test + "\t";
|
||||
out << setw(attW) << r.atts + "\t";
|
||||
out << setw(fieldW) << r.unit + "\t";
|
||||
if (r.GetMedian() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMedian() << "\t";
|
||||
if (r.GetMean() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMean() << "\t";
|
||||
if (r.GetStdDev() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetStdDev() << "\t";
|
||||
if (r.GetMin() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMin() << "\t";
|
||||
if (r.GetMax() == FLT_MAX)
|
||||
out << "N/A\t";
|
||||
else
|
||||
out << r.GetMax() << "\t";
|
||||
|
||||
out << endl;
|
||||
}
|
||||
out << endl
|
||||
<< "Note: results marked with (*) had missing values such as" << endl
|
||||
<< "might occur with a mixture of architectural capabilities." << endl;
|
||||
}
|
||||
|
||||
// ****************************************************************************
|
||||
// Method: ResultDatabase::ClearAllResults
|
||||
//
|
||||
// Purpose:
|
||||
// Clears all existing results from the ResultDatabase; used for multiple passes
|
||||
// of the same test or multiple tests.
|
||||
//
|
||||
// Arguments:
|
||||
//
|
||||
// Programmer: Jeffrey Young
|
||||
// Creation: September 10th, 2014
|
||||
//
|
||||
// Modifications:
|
||||
//
|
||||
//
|
||||
// ****************************************************************************
|
||||
void ResultDatabase::ClearAllResults() { results.clear(); }
|
||||
|
||||
// ****************************************************************************
|
||||
// Method: ResultDatabase::DumpCsv
|
||||
//
|
||||
// Purpose:
|
||||
// Writes either detailed or summary results (min/max/stddev/med/mean), but not
|
||||
// every individual trial.
|
||||
//
|
||||
// Arguments:
|
||||
// out file to print CSV results
|
||||
//
|
||||
// Programmer: Jeffrey Young
|
||||
// Creation: August 28th, 2014
|
||||
//
|
||||
// Modifications:
|
||||
//
|
||||
// ****************************************************************************
|
||||
void ResultDatabase::DumpCsv(string fileName) {
|
||||
bool emptyFile;
|
||||
vector<Result> sorted(results);
|
||||
|
||||
sort(sorted.begin(), sorted.end());
|
||||
|
||||
// Check to see if the file is empty - if so, add the headers
|
||||
emptyFile = this->IsFileEmpty(fileName);
|
||||
|
||||
// Open file and append by default
|
||||
ofstream out;
|
||||
out.open(fileName.c_str(), std::ofstream::out | std::ofstream::app);
|
||||
|
||||
// Add headers only for empty files
|
||||
if (emptyFile) {
|
||||
// TODO: in big parallel runs, the "trials" are the procs
|
||||
// and we really don't want to print them all out....
|
||||
out << "test, "
|
||||
<< "atts, "
|
||||
<< "units, "
|
||||
<< "median, "
|
||||
<< "mean, "
|
||||
<< "stddev, "
|
||||
<< "min, "
|
||||
<< "max, ";
|
||||
out << endl;
|
||||
}
|
||||
|
||||
for (int i = 0; i < sorted.size(); i++) {
|
||||
Result& r = sorted[i];
|
||||
out << r.test << ", ";
|
||||
out << r.atts << ", ";
|
||||
out << r.unit << ", ";
|
||||
if (r.GetMedian() == FLT_MAX)
|
||||
out << "N/A, ";
|
||||
else
|
||||
out << r.GetMedian() << ", ";
|
||||
if (r.GetMean() == FLT_MAX)
|
||||
out << "N/A, ";
|
||||
else
|
||||
out << r.GetMean() << ", ";
|
||||
if (r.GetStdDev() == FLT_MAX)
|
||||
out << "N/A, ";
|
||||
else
|
||||
out << r.GetStdDev() << ", ";
|
||||
if (r.GetMin() == FLT_MAX)
|
||||
out << "N/A, ";
|
||||
else
|
||||
out << r.GetMin() << ", ";
|
||||
if (r.GetMax() == FLT_MAX)
|
||||
out << "N/A, ";
|
||||
else
|
||||
out << r.GetMax() << ", ";
|
||||
|
||||
out << endl;
|
||||
}
|
||||
out << endl;
|
||||
|
||||
out.close();
|
||||
}
|
||||
|
||||
// ****************************************************************************
|
||||
// Method: ResultDatabase::IsFileEmpty
|
||||
//
|
||||
// Purpose:
|
||||
// Returns whether a file is empty - used as a helper for CSV printing
|
||||
//
|
||||
// Arguments:
|
||||
// file The input file to check for emptiness
|
||||
//
|
||||
// Programmer: Jeffrey Young
|
||||
// Creation: August 28th, 2014
|
||||
//
|
||||
// Modifications:
|
||||
//
|
||||
// ****************************************************************************
|
||||
|
||||
bool ResultDatabase::IsFileEmpty(string fileName) {
|
||||
|
||||
ifstream file(fileName.c_str());
|
||||
|
||||
// If the file doesn't exist it is by definition empty
|
||||
if (!file.good()) {
|
||||
return true;
|
||||
} else {
|
||||
bool fileEmpty;
|
||||
fileEmpty = (bool)(file.peek() == ifstream::traits_type::eof());
|
||||
file.close();
|
||||
|
||||
return fileEmpty;
|
||||
}
|
||||
|
||||
// Otherwise, return false
|
||||
return false;
|
||||
}
|
||||
|
||||
|
||||
// ****************************************************************************
|
||||
// Method: ResultDatabase::GetResultsForTest
|
||||
//
|
||||
// Purpose:
|
||||
// Returns a vector of results for just one test name.
|
||||
//
|
||||
// Arguments:
|
||||
// test the name of the test results to search for
|
||||
//
|
||||
// Programmer: Jeremy Meredith
|
||||
// Creation: December 3, 2010
|
||||
//
|
||||
// Modifications:
|
||||
//
|
||||
// ****************************************************************************
|
||||
vector<ResultDatabase::Result> ResultDatabase::GetResultsForTest(const string& test) {
|
||||
// get only the given test results
|
||||
vector<Result> retval;
|
||||
for (int i = 0; i < results.size(); i++) {
|
||||
Result& r = results[i];
|
||||
if (r.test == test) retval.push_back(r);
|
||||
}
|
||||
return retval;
|
||||
}
|
||||
|
||||
// ****************************************************************************
|
||||
// Method: ResultDatabase::GetResults
|
||||
//
|
||||
// Purpose:
|
||||
// Returns all the results.
|
||||
//
|
||||
// Arguments:
|
||||
//
|
||||
// Programmer: Jeremy Meredith
|
||||
// Creation: December 3, 2010
|
||||
//
|
||||
// Modifications:
|
||||
//
|
||||
// ****************************************************************************
|
||||
const vector<ResultDatabase::Result>& ResultDatabase::GetResults() const { return results; }
|
||||
@@ -1,89 +0,0 @@
|
||||
#ifndef RESULT_DATABASE_H
|
||||
#define RESULT_DATABASE_H
|
||||
|
||||
#include <string>
|
||||
#include <vector>
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <cfloat>
|
||||
using std::ifstream;
|
||||
using std::ofstream;
|
||||
using std::ostream;
|
||||
using std::string;
|
||||
using std::vector;
|
||||
|
||||
|
||||
// ****************************************************************************
|
||||
// Class: ResultDatabase
|
||||
//
|
||||
// Purpose:
|
||||
// Track numerical results as they are generated.
|
||||
// Print statistics of raw results.
|
||||
//
|
||||
// Programmer: Jeremy Meredith
|
||||
// Creation: June 12, 2009
|
||||
//
|
||||
// Modifications:
|
||||
// Jeremy Meredith, Wed Nov 10 14:20:47 EST 2010
|
||||
// Split timing reports into detailed and summary. E.g. for serial code,
|
||||
// we might report all trial values, but skip them in parallel.
|
||||
//
|
||||
// Jeremy Meredith, Thu Nov 11 11:40:18 EST 2010
|
||||
// Added check for missing value tag.
|
||||
//
|
||||
// Jeremy Meredith, Mon Nov 22 13:37:10 EST 2010
|
||||
// Added percentile statistic.
|
||||
//
|
||||
// Jeremy Meredith, Fri Dec 3 16:30:31 EST 2010
|
||||
// Added a method to extract a subset of results based on test name. Also,
|
||||
// the Result class is now public, so that clients can use them directly.
|
||||
// Added a GetResults method as well, and made several functions const.
|
||||
//
|
||||
// ****************************************************************************
|
||||
class ResultDatabase {
|
||||
public:
|
||||
//
|
||||
// A performance result for a single SHOC benchmark run.
|
||||
//
|
||||
struct Result {
|
||||
string test; // e.g. "readback"
|
||||
string atts; // e.g. "pagelocked 4k^2"
|
||||
string unit; // e.g. "MB/sec"
|
||||
vector<double> value; // e.g. "837.14"
|
||||
double GetMin() const;
|
||||
double GetMax() const;
|
||||
double GetMedian() const;
|
||||
double GetPercentile(double q) const;
|
||||
double GetMean() const;
|
||||
double GetStdDev() const;
|
||||
|
||||
bool operator<(const Result& rhs) const;
|
||||
|
||||
bool HadAnyFLTMAXValues() const {
|
||||
for (int i = 0; i < value.size(); ++i) {
|
||||
if (value[i] >= FLT_MAX) return true;
|
||||
}
|
||||
return false;
|
||||
}
|
||||
};
|
||||
|
||||
protected:
|
||||
vector<Result> results;
|
||||
|
||||
public:
|
||||
void AddResult(const string& test, const string& atts, const string& unit, double value);
|
||||
void AddResults(const string& test, const string& atts, const string& unit,
|
||||
const vector<double>& values);
|
||||
vector<Result> GetResultsForTest(const string& test);
|
||||
const vector<Result>& GetResults() const;
|
||||
void ClearAllResults();
|
||||
void DumpDetailed(ostream&);
|
||||
void DumpSummary(ostream&);
|
||||
void DumpCsv(string fileName);
|
||||
|
||||
private:
|
||||
bool IsFileEmpty(string fileName);
|
||||
};
|
||||
|
||||
|
||||
#endif
|
||||
@@ -1,3 +0,0 @@
|
||||
loop(1000); H2D; NullKernel; D2H; endloop;
|
||||
streamsync;
|
||||
printTiming(1000)
|
||||
@@ -1 +0,0 @@
|
||||
H2D; NullKernel; D2H; streamsync
|
||||
@@ -1,865 +0,0 @@
|
||||
#include <stdio.h>
|
||||
#include <iostream>
|
||||
#include <sstream>
|
||||
#include <iomanip>
|
||||
#include <algorithm>
|
||||
#include <string>
|
||||
#include <typeinfo>
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#ifndef _WIN32
|
||||
#include <sys/time.h>
|
||||
#endif
|
||||
|
||||
#include "ResultDatabase.h"
|
||||
#include "nullkernel.hip.cpp"
|
||||
|
||||
bool g_printedTiming = false;
|
||||
|
||||
// Cmdline parms:
|
||||
int p_device = 0;
|
||||
const char* p_command = "setstream(1); H2D; NullKernel; D2H;";
|
||||
const char* p_file = nullptr;
|
||||
unsigned p_verbose = 0x0;
|
||||
unsigned p_db = 0x0;
|
||||
unsigned p_blockingSync = 0x0;
|
||||
|
||||
//---
|
||||
int p_iterations = 1;
|
||||
|
||||
#define KNRM "\x1B[0m"
|
||||
#define KRED "\x1B[31m"
|
||||
#define KGRN "\x1B[32m"
|
||||
|
||||
|
||||
#define failed(...) \
|
||||
printf("error: "); \
|
||||
printf(__VA_ARGS__); \
|
||||
printf("\n"); \
|
||||
abort();
|
||||
|
||||
|
||||
#define HIPCHECK(error) \
|
||||
{ \
|
||||
hipError_t localError = error; \
|
||||
if (localError != hipSuccess) { \
|
||||
printf("%serror: '%s'(%d) from %s at %s:%d%s\n", KRED, hipGetErrorString(localError), \
|
||||
localError, #error, __FILE__, __LINE__, KNRM); \
|
||||
failed("API returned error code."); \
|
||||
} \
|
||||
}
|
||||
#define HIPASSERT(condition, msg) \
|
||||
if (!(condition)) { \
|
||||
failed("%sassertion %s at %s:%d: %s%s\n", KRED, #condition, __FILE__, __LINE__, msg, \
|
||||
KNRM); \
|
||||
}
|
||||
|
||||
|
||||
int parseInt(const char* str, int* output) {
|
||||
char* next;
|
||||
*output = strtol(str, &next, 0);
|
||||
return !strlen(next);
|
||||
}
|
||||
|
||||
|
||||
void printConfig() {
|
||||
hipDeviceProp_t props;
|
||||
HIPCHECK(hipGetDeviceProperties(&props, p_device));
|
||||
|
||||
printf("Device:%s Mem=%.1fGB #CUs=%d Freq=%.0fMhz\n", props.name,
|
||||
props.totalGlobalMem / 1024.0 / 1024.0 / 1024.0, props.multiProcessorCount,
|
||||
props.clockRate / 1000.0);
|
||||
}
|
||||
|
||||
|
||||
void help() {
|
||||
printf("Usage: hipBusBandwidth [OPTIONS]\n");
|
||||
printf(" --file, -f : Read string of commands from file\n");
|
||||
printf(" --command, -c : String specifying commands to run.\n");
|
||||
printf(" --iterations, -i : Number of copy iterations to run.\n");
|
||||
printf(" --device, -d : Device ID to use (0..numDevices).\n");
|
||||
printf(
|
||||
" --verbose, -v : Verbose printing of status. Fore more info, combine with "
|
||||
"HIP_TRACE_API on ROCm\n");
|
||||
};
|
||||
|
||||
|
||||
int parseStandardArguments(int argc, char* argv[]) {
|
||||
for (int i = 1; i < argc; i++) {
|
||||
const char* arg = argv[i];
|
||||
|
||||
if (!strcmp(arg, " ")) {
|
||||
// skip NULL args.
|
||||
} else if (!strcmp(arg, "--iterations") || (!strcmp(arg, "-i"))) {
|
||||
if (++i >= argc || !parseInt(argv[i], &p_iterations)) {
|
||||
failed("Bad --iterations argument");
|
||||
}
|
||||
|
||||
} else if (!strcmp(arg, "--device") || (!strcmp(arg, "-d"))) {
|
||||
if (++i >= argc || !parseInt(argv[i], &p_device)) {
|
||||
failed("Bad --device argument");
|
||||
}
|
||||
|
||||
} else if (!strcmp(arg, "--file") || (!strcmp(arg, "-f"))) {
|
||||
if (++i >= argc) {
|
||||
failed("Bad --file argument");
|
||||
} else {
|
||||
p_file = argv[i];
|
||||
}
|
||||
|
||||
} else if (!strcmp(arg, "--commands") || (!strcmp(arg, "-c"))) {
|
||||
if (++i >= argc) {
|
||||
failed("Bad --commands argument");
|
||||
} else {
|
||||
p_command = argv[i];
|
||||
}
|
||||
|
||||
} else if (!strcmp(arg, "--verbose") || (!strcmp(arg, "-v"))) {
|
||||
p_verbose = 1;
|
||||
|
||||
} else if (!strcmp(arg, "--blockingSync") || (!strcmp(arg, "-B"))) {
|
||||
p_blockingSync = 1;
|
||||
|
||||
|
||||
} else if (!strcmp(arg, "--help") || (!strcmp(arg, "-h"))) {
|
||||
help();
|
||||
exit(EXIT_SUCCESS);
|
||||
|
||||
} else {
|
||||
failed("Bad argument '%s'", arg);
|
||||
}
|
||||
}
|
||||
|
||||
return 0;
|
||||
};
|
||||
|
||||
// Returns the current system time in microseconds
|
||||
inline long long get_time() {
|
||||
#ifdef _WIN32
|
||||
struct timespec ts;
|
||||
timespec_get(&ts, TIME_UTC);
|
||||
return (ts.tv_sec * 1000000) + (ts.tv_nsec/1000);
|
||||
#else
|
||||
struct timeval tv;
|
||||
gettimeofday(&tv, 0);
|
||||
return (tv.tv_sec * 1000000) + tv.tv_usec;
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
class Command;
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
// A stream of commands , specified as a string.
|
||||
class CommandStream {
|
||||
public:
|
||||
// State that is inherited by sub-blocks:
|
||||
struct CommandStreamState {
|
||||
hipStream_t _currentStream;
|
||||
std::vector<hipStream_t> _streams;
|
||||
vector<CommandStream*> _subBlocks;
|
||||
};
|
||||
|
||||
public:
|
||||
CommandStream(std::string commandStreamString, int iterations);
|
||||
~CommandStream();
|
||||
|
||||
hipStream_t currentStream() const { return _state._currentStream; };
|
||||
|
||||
void print(const std::string& indent = "") const;
|
||||
void printBrief(std::ostream& s = std::cout) const;
|
||||
void run();
|
||||
void recordTime();
|
||||
void printTiming(int iterations = 0);
|
||||
|
||||
CommandStream* currentCommandStream() {
|
||||
return _parseInSubBlock ? _state._subBlocks.back() : this;
|
||||
};
|
||||
|
||||
void enterSubBlock(CommandStream* commandStream) {
|
||||
_parseInSubBlock = true;
|
||||
_state._subBlocks.push_back(commandStream);
|
||||
};
|
||||
|
||||
void exitSubBlock() { _parseInSubBlock = false; };
|
||||
|
||||
|
||||
void setParent(CommandStream* parentCmdStream) {
|
||||
_parentCommandStream = parentCmdStream;
|
||||
_state = parentCmdStream->_state;
|
||||
};
|
||||
CommandStream* getParent() { return _parentCommandStream; };
|
||||
|
||||
void setStream(int streamIndex);
|
||||
|
||||
CommandStreamState& getState() { return _state; };
|
||||
|
||||
private:
|
||||
static void tokenize(const std::string& s, char delim, std::vector<std::string>& tokens);
|
||||
void parse(const std::string fullCmd);
|
||||
|
||||
protected:
|
||||
CommandStreamState _state;
|
||||
|
||||
private:
|
||||
// List of commands to run in this stream:
|
||||
std::vector<Command*> _commands;
|
||||
|
||||
|
||||
// Number of iterations to run the command loop
|
||||
int _iterations;
|
||||
|
||||
|
||||
// Us to run the the command-stream. Only valid after run is called.
|
||||
long long _startTime;
|
||||
double _elapsedUs;
|
||||
|
||||
// Track nested loop of command streams:
|
||||
CommandStream* _parentCommandStream;
|
||||
|
||||
// Track if we are parsing commands in the subblock.
|
||||
bool _parseInSubBlock;
|
||||
};
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
class Command {
|
||||
public:
|
||||
// @p minArgs : Minimum arguments for command. -1 = don't check.
|
||||
// @p maxArgs : Minimum arguments for command. 0 means min=max, ie exact #arguments expected.
|
||||
// -1 = don't check max.
|
||||
Command(CommandStream* cmdStream, const std::vector<std::string>& args, int minArgs = 0,
|
||||
int maxArgs = 0)
|
||||
: _commandStream(cmdStream), _args(args) {
|
||||
int numArgs = args.size() - 1;
|
||||
|
||||
if ((minArgs != -1) && (numArgs < minArgs)) {
|
||||
// TODO - print full command here.
|
||||
failed("Not enough arguments for command %s. (Expected %d, got %d)", args[0].c_str(),
|
||||
minArgs, numArgs);
|
||||
}
|
||||
|
||||
// Check for an exact number of arguments:
|
||||
if (maxArgs == 0) {
|
||||
maxArgs = minArgs;
|
||||
}
|
||||
if ((maxArgs != -1) && (numArgs > maxArgs)) {
|
||||
failed("Too many arguments for command %s. (Expected %d, got %d)", args[0].c_str(),
|
||||
maxArgs, numArgs);
|
||||
}
|
||||
};
|
||||
|
||||
void printBrief(std::ostream& s = std::cout) const { s << _args[0]; }
|
||||
|
||||
virtual ~Command(){};
|
||||
|
||||
virtual void print(const std::string& indent = "") const {
|
||||
std::cout << indent << "[";
|
||||
std::for_each(_args.begin(), _args.end(), [](const std::string& s) { std::cout << s; });
|
||||
std::cout << "]";
|
||||
};
|
||||
|
||||
|
||||
virtual void run() = 0;
|
||||
|
||||
protected:
|
||||
int readIntArg(int argIndex, const std::string& argName) {
|
||||
// TODO - catch references to non-existant arguments here.
|
||||
int argVal;
|
||||
try {
|
||||
argVal = std::stoi(_args[argIndex]);
|
||||
} catch (std::invalid_argument) {
|
||||
failed("Command %s has bad %s argument ('%s')", _args[0].c_str(), argName.c_str(),
|
||||
_args[argIndex].c_str());
|
||||
}
|
||||
return argVal;
|
||||
}
|
||||
|
||||
protected:
|
||||
CommandStream* _commandStream;
|
||||
std::vector<std::string> _args;
|
||||
};
|
||||
|
||||
|
||||
#define FILENAME "nullkernel.hsaco"
|
||||
#define KERNEL_NAME "NullKernel"
|
||||
|
||||
//=================================================================================================
|
||||
// HCC optimizes away fully NULL kernel calls, so run one that is nearly null:
|
||||
class ModuleKernelCommand : public Command {
|
||||
public:
|
||||
ModuleKernelCommand(CommandStream* cmdStream, const std::vector<std::string>& args)
|
||||
: Command(cmdStream, args), _stream(cmdStream->currentStream()) {
|
||||
hipModule_t module;
|
||||
HIPCHECK(hipModuleLoad(&module, FILENAME));
|
||||
HIPCHECK(hipModuleGetFunction(&_function, module, KERNEL_NAME));
|
||||
};
|
||||
~ModuleKernelCommand(){};
|
||||
|
||||
void run() override {
|
||||
#define LEN 64
|
||||
float *X = NULL;
|
||||
HIPCHECK(hipMalloc((void**)&X, sizeof(float)));
|
||||
struct {
|
||||
float *Ad;
|
||||
}args;
|
||||
args.Ad = X;
|
||||
size_t argSize = sizeof(args);
|
||||
|
||||
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE, &argSize, HIP_LAUNCH_PARAM_END};
|
||||
|
||||
hipModuleLaunchKernel(_function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config);
|
||||
};
|
||||
|
||||
|
||||
public:
|
||||
hipFunction_t _function;
|
||||
hipStream_t _stream;
|
||||
};
|
||||
|
||||
|
||||
class KernelCommand : public Command {
|
||||
public:
|
||||
enum Type { Null, VectorAdd };
|
||||
KernelCommand(CommandStream* cmdStream, const std::vector<std::string>& args, Type kind)
|
||||
: Command(cmdStream, args), _kind(kind), _stream(cmdStream->currentStream()){};
|
||||
~KernelCommand(){};
|
||||
|
||||
|
||||
void run() override {
|
||||
static const int gridX = 64;
|
||||
static const int groupX = 64;
|
||||
|
||||
switch (_kind) {
|
||||
case Null:
|
||||
hipLaunchKernelGGL(NullKernel, dim3(gridX / groupX), dim3(gridX), 0, _stream, nullptr);
|
||||
break;
|
||||
case VectorAdd:
|
||||
assert(0); // TODO
|
||||
break;
|
||||
};
|
||||
}
|
||||
|
||||
private:
|
||||
Type _kind;
|
||||
hipStream_t _stream;
|
||||
};
|
||||
|
||||
//=================================================================================================
|
||||
class CopyCommand : public Command {
|
||||
enum MemType { PinnedHost, UnpinnedHost, Device };
|
||||
|
||||
public:
|
||||
CopyCommand(CommandStream* cmdStream, const std::vector<std::string>& args, hipMemcpyKind kind,
|
||||
bool isAsync, bool isPinnedHost);
|
||||
|
||||
~CopyCommand() {
|
||||
if (_dst) {
|
||||
dealloc(_dst, _dstType);
|
||||
_dst = NULL;
|
||||
};
|
||||
|
||||
if (_src) {
|
||||
dealloc(_src, _srcType);
|
||||
_src = NULL;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
void run() override {
|
||||
if (_isAsync) {
|
||||
HIPCHECK(hipMemcpyAsync(_dst, _src, _sizeBytes, _kind, _stream));
|
||||
} else {
|
||||
HIPCHECK(hipMemcpy(_dst, _src, _sizeBytes, _kind));
|
||||
}
|
||||
};
|
||||
|
||||
private:
|
||||
void* alloc(size_t size, MemType memType) {
|
||||
void* p;
|
||||
if (memType == Device) {
|
||||
HIPCHECK(hipMalloc(&p, size));
|
||||
|
||||
} else if (memType == PinnedHost) {
|
||||
HIPCHECK(hipHostMalloc(&p, size));
|
||||
|
||||
} else if (memType == UnpinnedHost) {
|
||||
p = (char*)malloc(size);
|
||||
HIPASSERT(p, "malloc failed");
|
||||
|
||||
} else {
|
||||
HIPASSERT(0, "unsupported memType");
|
||||
}
|
||||
|
||||
return p;
|
||||
};
|
||||
|
||||
|
||||
static void dealloc(void* p, MemType memType) {
|
||||
if (memType == Device) {
|
||||
HIPCHECK(hipFree(p));
|
||||
} else if (memType == PinnedHost) {
|
||||
HIPCHECK(hipHostFree(p));
|
||||
} else if (memType == UnpinnedHost) {
|
||||
free(p);
|
||||
} else {
|
||||
HIPASSERT(0, "unsupported memType");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
private:
|
||||
bool _isAsync;
|
||||
hipStream_t _stream;
|
||||
hipMemcpyKind _kind;
|
||||
|
||||
size_t _sizeBytes;
|
||||
void* _dst;
|
||||
MemType _dstType;
|
||||
|
||||
void* _src;
|
||||
MemType _srcType;
|
||||
};
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
class DeviceSyncCommand : public Command {
|
||||
public:
|
||||
DeviceSyncCommand(CommandStream* cmdStream, const std::vector<std::string>& args)
|
||||
: Command(cmdStream, args){};
|
||||
|
||||
void run() override { HIPCHECK(hipDeviceSynchronize()); };
|
||||
};
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
class StreamSyncCommand : public Command {
|
||||
public:
|
||||
StreamSyncCommand(CommandStream* cmdStream, const std::vector<std::string>& args)
|
||||
: Command(cmdStream, args), _stream(cmdStream->currentStream()){};
|
||||
|
||||
static const char* help() { return "synchronizes the current stream"; };
|
||||
|
||||
|
||||
void run() override { HIPCHECK(hipStreamSynchronize(_stream)); };
|
||||
|
||||
private:
|
||||
hipStream_t _stream;
|
||||
};
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
|
||||
//=================================================================================================
|
||||
class LoopCommand : public Command {
|
||||
public:
|
||||
LoopCommand(CommandStream* parentCmdStream, const std::vector<std::string>& args)
|
||||
: Command(parentCmdStream, args, 1) {
|
||||
int loopCnt;
|
||||
try {
|
||||
loopCnt = std::stoi(args[1]);
|
||||
} catch (std::invalid_argument) {
|
||||
failed("bad LOOP_CNT=%s", args[1].c_str());
|
||||
}
|
||||
|
||||
_commandStream = new CommandStream("", loopCnt);
|
||||
_commandStream->setParent(parentCmdStream);
|
||||
parentCmdStream->enterSubBlock(_commandStream);
|
||||
};
|
||||
|
||||
|
||||
void print(const std::string& indent = "") const override {
|
||||
Command::print();
|
||||
_commandStream->print(indent + " ");
|
||||
};
|
||||
|
||||
void run() override { _commandStream->run(); };
|
||||
};
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
class EndBlockCommand : public Command {
|
||||
public:
|
||||
EndBlockCommand(CommandStream* blockCmdStream, CommandStream* parentCmdStream,
|
||||
const std::vector<std::string>& args)
|
||||
: Command(parentCmdStream, args, 0, 1), _blockCmdStream(blockCmdStream), _printTiming(0) {
|
||||
int argCnt = args.size() - 1;
|
||||
if (argCnt >= 1) {
|
||||
_printTiming = readIntArg(1, "PRINT_TIMING");
|
||||
}
|
||||
|
||||
if (parentCmdStream == nullptr) {
|
||||
failed("%s without corresponding command to start block", args[0].c_str());
|
||||
}
|
||||
parentCmdStream->exitSubBlock();
|
||||
};
|
||||
|
||||
void run() override {
|
||||
if (_printTiming) {
|
||||
_blockCmdStream->printTiming();
|
||||
}
|
||||
};
|
||||
|
||||
private:
|
||||
CommandStream* _blockCmdStream;
|
||||
|
||||
// print the stream when loop exits.
|
||||
int _printTiming;
|
||||
};
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
class SetStreamCommand : public Command {
|
||||
public:
|
||||
SetStreamCommand(CommandStream* cmdStream, const std::vector<std::string>& args)
|
||||
: Command(cmdStream, args, 1) {
|
||||
int streamIndex = readIntArg(1, "STREAM_INDEX");
|
||||
|
||||
cmdStream->setStream(streamIndex);
|
||||
};
|
||||
|
||||
void run() override{};
|
||||
};
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
class PrintTimingCommand : public Command {
|
||||
public:
|
||||
PrintTimingCommand(CommandStream* cmdStream, const std::vector<std::string>& args)
|
||||
: Command(cmdStream, args, 1) {
|
||||
_iterations = readIntArg(1, "ITERATIONS");
|
||||
};
|
||||
|
||||
void run() override { _commandStream->printTiming(_iterations); };
|
||||
|
||||
private:
|
||||
int _iterations;
|
||||
};
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
CopyCommand::CopyCommand(CommandStream* cmdStream, const std::vector<std::string>& args,
|
||||
hipMemcpyKind kind, bool isAsync, bool isPinnedHost)
|
||||
: Command(cmdStream, args),
|
||||
_isAsync(isAsync),
|
||||
_stream(cmdStream->currentStream()),
|
||||
_kind(kind) {
|
||||
switch (kind) {
|
||||
case hipMemcpyDeviceToHost:
|
||||
_srcType = Device;
|
||||
_dstType = isPinnedHost ? PinnedHost : UnpinnedHost;
|
||||
break;
|
||||
case hipMemcpyHostToDevice:
|
||||
_srcType = isPinnedHost ? PinnedHost : UnpinnedHost;
|
||||
_dstType = Device;
|
||||
break;
|
||||
default:
|
||||
HIPASSERT(0, "Unknown hipMemcpyKind");
|
||||
};
|
||||
|
||||
_sizeBytes = 64; // TODO, support reading from arg.
|
||||
|
||||
_dst = alloc(_sizeBytes, _dstType);
|
||||
_src = alloc(_sizeBytes, _srcType);
|
||||
};
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
//=================================================================================================
|
||||
// Implementations:
|
||||
//=================================================================================================
|
||||
|
||||
//=================================================================================================
|
||||
CommandStream::CommandStream(std::string commandStreamString, int iterations)
|
||||
: _iterations(iterations),
|
||||
_startTime(0),
|
||||
_elapsedUs(0.0),
|
||||
_parentCommandStream(nullptr),
|
||||
_parseInSubBlock(false) {
|
||||
std::vector<std::string> tokens;
|
||||
tokenize(commandStreamString, ';', tokens);
|
||||
|
||||
setStream(0);
|
||||
std::for_each(tokens.begin(), tokens.end(), [&](const std::string s) { this->parse(s); });
|
||||
}
|
||||
|
||||
|
||||
CommandStream::~CommandStream() {
|
||||
std::for_each(_state._streams.begin(), _state._streams.end(), [&](hipStream_t s) {
|
||||
if (s) {
|
||||
HIPCHECK(hipStreamDestroy(s));
|
||||
}
|
||||
});
|
||||
|
||||
std::for_each(_commands.begin(), _commands.end(), [&](Command* c) { delete c; });
|
||||
}
|
||||
|
||||
|
||||
void CommandStream::setStream(int streamIndex) {
|
||||
if (streamIndex >= _state._streams.size()) {
|
||||
_state._streams.resize(streamIndex + 1);
|
||||
}
|
||||
|
||||
if (streamIndex && (_state._streams[streamIndex] == nullptr)) {
|
||||
// Create new stream:
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
_state._streams[streamIndex] = stream;
|
||||
_state._currentStream = stream;
|
||||
} else {
|
||||
// Use existing stream:
|
||||
|
||||
_state._currentStream = _state._streams[streamIndex];
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
void CommandStream::tokenize(const std::string& s, char delim, std::vector<std::string>& tokens) {
|
||||
std::stringstream ss;
|
||||
ss.str(s);
|
||||
std::string item;
|
||||
while (getline(ss, item, delim)) {
|
||||
item.erase(std::remove(item.begin(), item.end(), ' '), item.end()); // remove whitespace.
|
||||
tokens.push_back(item);
|
||||
}
|
||||
}
|
||||
|
||||
void trim(std::string* s) {
|
||||
// trim whitespace from begin and end:
|
||||
const char* t = "\t\n\r\f\v";
|
||||
s->erase(0, s->find_first_not_of(t));
|
||||
s->erase(s->find_last_not_of(t) + 1);
|
||||
}
|
||||
|
||||
void ltrim(std::string* s) {
|
||||
// trim whitespace from begin and end:
|
||||
const char* t = "\t\n\r\f\v";
|
||||
s->erase(0, s->find_first_not_of(t));
|
||||
}
|
||||
|
||||
void CommandStream::parse(std::string fullCmd) {
|
||||
// convert to lower-case:
|
||||
std::transform(fullCmd.begin(), fullCmd.end(), fullCmd.begin(), ::tolower);
|
||||
trim(&fullCmd);
|
||||
|
||||
if (p_db) {
|
||||
printf("parse: <%s>\n", fullCmd.c_str());
|
||||
}
|
||||
|
||||
|
||||
std::string c;
|
||||
std::vector<std::string> args;
|
||||
size_t leftParenZ = fullCmd.find_first_of('(');
|
||||
if (leftParenZ == string::npos) {
|
||||
c = fullCmd;
|
||||
args.push_back(c);
|
||||
} else {
|
||||
c = fullCmd.substr(0, leftParenZ);
|
||||
args.push_back(c);
|
||||
size_t rightParenZ = fullCmd.find_first_of(')', leftParenZ);
|
||||
std::string argStr = fullCmd.substr(leftParenZ + 1, rightParenZ - leftParenZ - 1);
|
||||
// printf ("c=%s argstr='%s' leftParenZ=%zu rightParenZ=%zu\n", c.c_str(), argStr.c_str(),
|
||||
// leftParenZ, rightParenZ);
|
||||
tokenize(argStr, ',', args);
|
||||
}
|
||||
|
||||
|
||||
if ((args.size() == 0) || (fullCmd.c_str()[0] == '#')) {
|
||||
if (p_db) {
|
||||
printf(" skip comment\n");
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
|
||||
Command* cmd = NULL;
|
||||
CommandStream* cmdStream = currentCommandStream();
|
||||
|
||||
if (c == "h2d") {
|
||||
cmd = new CopyCommand(cmdStream, args, hipMemcpyHostToDevice, true /*isAsync*/,
|
||||
true /*isPinned*/);
|
||||
//= h2d
|
||||
//= Performs an async host-to-device copy of array A_h to A_d.
|
||||
//= The size of these arrays may be set with the datasize command.
|
||||
|
||||
} else if (c == "d2h") {
|
||||
cmd = new CopyCommand(cmdStream, args, hipMemcpyDeviceToHost, true /*isAsync*/,
|
||||
true /*isPinned*/);
|
||||
//= d2h
|
||||
//= Performs an async device-to-host copy of array A_d to A_h.
|
||||
//= The size of these arrays may be set with the datasize command.
|
||||
|
||||
} else if (c == "modulekernel") {
|
||||
cmd = new ModuleKernelCommand(cmdStream, args);
|
||||
|
||||
} else if (c == "nullkernel") {
|
||||
cmd = new KernelCommand(cmdStream, args, KernelCommand::Null);
|
||||
//= nullkernel
|
||||
//= Dispatches a null kernel to the device.
|
||||
|
||||
} else if (c == "vectoraddkernel") {
|
||||
cmd = new KernelCommand(cmdStream, args, KernelCommand::VectorAdd);
|
||||
|
||||
} else if (c == "devicesync") {
|
||||
cmd = new DeviceSyncCommand(cmdStream, args);
|
||||
|
||||
} else if (c == "streamsync") {
|
||||
//= streamsync
|
||||
//= Execute hipStreamSynchronize.
|
||||
//= This will cause the host thread to wait until the current stream
|
||||
//= completes all pending operations.
|
||||
cmd = new StreamSyncCommand(cmdStream, args);
|
||||
|
||||
} else if (c == "setstream") {
|
||||
//= setstream(STREAM_INDEX);
|
||||
//= Set current stream used by subsequent commands.
|
||||
//= STREAM_INDEX is index starting from 0...N.
|
||||
//= This function will create new stream on first call to setstream or re-use previous
|
||||
//= stream if setstream has already been called with STREAM_INDEX.
|
||||
//= STREAM_INDEX=0 will use the default "null" stream associated with the device, and will
|
||||
//not create a new stream. = The default stream has special, conservative synchronization
|
||||
//properties.
|
||||
|
||||
cmd = new SetStreamCommand(cmdStream, args);
|
||||
|
||||
} else if (c == "printtiming") {
|
||||
cmd = new PrintTimingCommand(cmdStream, args);
|
||||
|
||||
} else if (c == "loop") {
|
||||
//= loop(LOOP_CNT)
|
||||
//= Loop over next set of commands (until 'endloop' command) for LOOP_CNT iterations.
|
||||
//= Loops can be nested.
|
||||
|
||||
cmd = new LoopCommand(cmdStream, args);
|
||||
|
||||
} else if (c == "endloop") {
|
||||
//= endloop
|
||||
//= End a looped sequence. Must be paired with a preceding loop command.
|
||||
//= Command between the `loop` and `endloop` must be executed
|
||||
|
||||
CommandStream* parentCmdStream = cmdStream->getParent();
|
||||
cmd = new EndBlockCommand(cmdStream, parentCmdStream, args);
|
||||
cmdStream = parentCmdStream;
|
||||
|
||||
} else {
|
||||
std::cerr << "error: Bad command '" << fullCmd << "\n";
|
||||
HIPASSERT(0, "bad command in command-stream");
|
||||
}
|
||||
|
||||
if (cmd) {
|
||||
cmdStream->_commands.push_back(cmd);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
void CommandStream::print(const std::string& indent) const {
|
||||
for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) {
|
||||
(*cmdI)->print(indent);
|
||||
};
|
||||
}
|
||||
|
||||
|
||||
void CommandStream::printBrief(std::ostream& s) const {
|
||||
for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) {
|
||||
(*cmdI)->printBrief(s);
|
||||
s << ";";
|
||||
};
|
||||
}
|
||||
|
||||
void CommandStream::run() {
|
||||
_startTime = get_time();
|
||||
for (int i = 0; i < _iterations; i++) {
|
||||
for (auto cmdI = _commands.begin(); cmdI != _commands.end(); cmdI++) {
|
||||
if (p_verbose) {
|
||||
(*cmdI)->print();
|
||||
}
|
||||
(*cmdI)->run();
|
||||
}
|
||||
}
|
||||
|
||||
// Record time, if not already stored. (an earlier printTime command will also store the time)
|
||||
recordTime();
|
||||
};
|
||||
|
||||
void CommandStream::recordTime() {
|
||||
if (_elapsedUs == 0.0) {
|
||||
auto stopTime = get_time();
|
||||
_elapsedUs = stopTime - _startTime;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
void CommandStream::printTiming(int iterations) {
|
||||
if ((_state._subBlocks.size() == 1) && (_commands.size() == 1)) {
|
||||
// printf ("print just the loop\n");
|
||||
_state._subBlocks.front()->printTiming(iterations);
|
||||
} else {
|
||||
g_printedTiming = true;
|
||||
|
||||
recordTime();
|
||||
if (iterations == 0) {
|
||||
iterations = _iterations;
|
||||
}
|
||||
std::cout << "command<";
|
||||
printBrief(std::cout);
|
||||
std::cout << ">,";
|
||||
printf(" iterations,%d, total_time,%6.3f, time/iteration,%6.3f\n", iterations,
|
||||
_elapsedUs, _elapsedUs / iterations);
|
||||
}
|
||||
};
|
||||
|
||||
|
||||
//=================================================================================================
|
||||
int main(int argc, char* argv[]) {
|
||||
parseStandardArguments(argc, argv);
|
||||
|
||||
printConfig();
|
||||
|
||||
CommandStream* cs;
|
||||
|
||||
if (p_blockingSync) {
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
printf("setting BlockingSync for AMD\n");
|
||||
#ifdef _WIN32
|
||||
_putenv_s("HIP_BLOCKING_SYNC", "1");
|
||||
#else
|
||||
setenv("HIP_BLOCKING_SYNC", "1", 1);
|
||||
#endif
|
||||
#endif
|
||||
#ifdef __HIP_PLATFORM_NVIDIA__
|
||||
printf("setting cudaDeviceBlockingSync\n");
|
||||
HIPCHECK(hipSetDeviceFlags(cudaDeviceBlockingSync));
|
||||
#endif
|
||||
};
|
||||
|
||||
|
||||
if (p_file) {
|
||||
// TODO - catch exception on file IO here:
|
||||
std::ifstream file(p_file);
|
||||
std::string str;
|
||||
std::string file_contents;
|
||||
while (std::getline(file, str)) {
|
||||
file_contents += str;
|
||||
}
|
||||
|
||||
cs = new CommandStream(file_contents, p_iterations);
|
||||
|
||||
} else {
|
||||
cs = new CommandStream(p_command, p_iterations);
|
||||
}
|
||||
|
||||
cs->print();
|
||||
printf("------\n");
|
||||
|
||||
cs->run();
|
||||
if (!g_printedTiming) {
|
||||
cs->printTiming();
|
||||
}
|
||||
|
||||
delete cs;
|
||||
}
|
||||
|
||||
|
||||
// TODO - add error checking for arguments.
|
||||
@@ -1,3 +0,0 @@
|
||||
setstream(1);
|
||||
NullKernel; streamsync;
|
||||
loop(10000); H2D; NullKernel; streamsync; endloop(1);
|
||||
@@ -1,3 +0,0 @@
|
||||
loop(1000); H2D; NullKernel; D2H; endloop;
|
||||
streamsync;
|
||||
printTiming(1000)
|
||||
@@ -1,2 +0,0 @@
|
||||
setstream(1);
|
||||
loop(1000); NullKernel; streamsync; endloop(1);
|
||||
@@ -1,7 +0,0 @@
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
extern "C" __global__ void NullKernel(float* Ad) {
|
||||
if (Ad) {
|
||||
Ad[0] = 42;
|
||||
}
|
||||
}
|
||||
Binární soubor nebyl zobrazen.
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
NullKernel; streamsync;
|
||||
loop(30000); NullKernel; streamsync; endloop(1);
|
||||
loop(30000); H2D; H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(30000); H2D; H2D; H2D; NullKernel; streamsync; endloop(1);
|
||||
|
||||
loop(30000); H2D; NullKernel; D2H; streamsync; endloop(1);
|
||||
loop(30000); NullKernel; D2H; streamsync; endloop(1);
|
||||
loop(30000); NullKernel; D2H; D2H; streamsync; endloop(1);
|
||||
loop(30000); NullKernel; D2H; D2H; D2H; streamsync; endloop(1);
|
||||
@@ -1,8 +0,0 @@
|
||||
setstream(1);
|
||||
NullKernel; streamsync;
|
||||
loop(100000); NullKernel; streamsync; endloop(1);
|
||||
|
||||
loop(100000); H2D; streamsync; NullKernel; streamsync; endloop(1);
|
||||
|
||||
loop(100000); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
|
||||
@@ -1,5 +0,0 @@
|
||||
setstream(1);
|
||||
NullKernel; streamsync;
|
||||
loop(100000); NullKernel; streamsync; endloop(1);
|
||||
loop(100000); H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(100000); H2D; NullKernel; D2H; streamsync; endloop(1);
|
||||
@@ -1,7 +0,0 @@
|
||||
setstream(0);
|
||||
NullKernel; streamsync;
|
||||
loop(100000); NullKernel; streamsync; endloop(1);
|
||||
|
||||
loop(100000); H2D; NullKernel; streamsync; endloop(1);
|
||||
|
||||
loop(100000); H2D; NullKernel; D2H; streamsync; endloop(1);
|
||||
@@ -1,5 +0,0 @@
|
||||
setstream(1);
|
||||
NullKernel; streamsync;
|
||||
loop(100); ModuleKernel; streamsync; endloop(1);
|
||||
loop(100); streamsync; endloop(1);
|
||||
loop(3000); NullKernel; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
|
||||
loop(10); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
|
||||
loop(100); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
|
||||
loop(100); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
|
||||
loop(1000); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
|
||||
loop(1000); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
|
||||
loop(10000); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
|
||||
loop(10000); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(10); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(100); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(100); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(10); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
|
||||
loop(10); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
|
||||
loop(100); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
|
||||
loop(100); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
|
||||
loop(1000); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
|
||||
loop(1000); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
|
||||
loop(10000); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
|
||||
loop(10000); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(10); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1);
|
||||
loop(100); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1);
|
||||
loop(100); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(100); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(100); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(1000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(1000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(100); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(100); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
|
||||
loop(10); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
|
||||
loop(100); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
|
||||
loop(100); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
|
||||
loop(10); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
|
||||
loop(100); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
|
||||
loop(100); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(100); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(100); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(100); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(100); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); streamsync; streamsync; endloop(1);
|
||||
loop(10); streamsync; streamsync; endloop(1);
|
||||
loop(100); streamsync; streamsync; endloop(1);
|
||||
loop(100); streamsync; streamsync; endloop(1);
|
||||
loop(1000); streamsync; streamsync; endloop(1);
|
||||
loop(1000); streamsync; streamsync; endloop(1);
|
||||
loop(10000); streamsync; streamsync; endloop(1);
|
||||
loop(10000); streamsync; streamsync; endloop(1);
|
||||
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10); D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(1000); D2H;streamsync; D2H; streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; D2H; streamsync; endloop(1);
|
||||
loop(10); D2H; D2H; streamsync; endloop(1);
|
||||
loop(100); D2H; D2H; streamsync; endloop(1);
|
||||
loop(100); D2H; D2H; streamsync; endloop(1);
|
||||
loop(1000); D2H; D2H; streamsync; endloop(1);
|
||||
loop(1000); D2H; D2H; streamsync; endloop(1);
|
||||
loop(1000); D2H; D2H; streamsync; endloop(1);
|
||||
loop(10000); D2H; D2H; streamsync; endloop(1);
|
||||
loop(10000); D2H; D2H; streamsync; endloop(1);
|
||||
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10); H2D; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(100); H2D; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(100); H2D; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D;streamsync; H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10000); H2D; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10000); H2D; streamsync; H2D; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1);
|
||||
loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(100); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1);
|
||||
loop(100); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1);
|
||||
loop(1000); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1);
|
||||
loop(1000); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1);
|
||||
loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(10); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1);
|
||||
loop(100); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1);
|
||||
loop(100); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1);
|
||||
loop(1000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); H2D ; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1);
|
||||
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; H2D; streamsync; endloop(1);
|
||||
loop(10); H2D; H2D; streamsync; endloop(1);
|
||||
loop(100); H2D; H2D; streamsync; endloop(1);
|
||||
loop(100); H2D; H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D; H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D; H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D; H2D; streamsync; endloop(1);
|
||||
loop(10000); H2D; H2D; streamsync; endloop(1);
|
||||
loop(10000); H2D; H2D; streamsync; endloop(1);
|
||||
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(10); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(100); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(100); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
-10
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(10); NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(100); NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(100); NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; NullKernel; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
|
||||
loop(10); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
|
||||
loop(100); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
|
||||
loop(100); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
|
||||
loop(1000); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
|
||||
loop(1000); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
|
||||
loop(10000); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
|
||||
loop(10000); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(10); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(100); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(100); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
|
||||
loop(10); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
|
||||
loop(10); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1);
|
||||
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
|
||||
loop(10); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
|
||||
loop(100); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
|
||||
loop(100); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
|
||||
loop(1000); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
|
||||
loop(1000); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
|
||||
loop(10000); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
|
||||
loop(10000); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1);
|
||||
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(10); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(100); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(100); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
|
||||
loop(10); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
|
||||
loop(100); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
|
||||
loop(100); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
|
||||
loop(1000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
|
||||
loop(1000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
|
||||
loop(10000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
|
||||
loop(10000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
|
||||
loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
|
||||
loop(100); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
|
||||
loop(100); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
|
||||
loop(1000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
|
||||
loop(1000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
|
||||
loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
|
||||
loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
|
||||
loop(10); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
|
||||
loop(100); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
|
||||
loop(100); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
|
||||
loop(1000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
|
||||
loop(1000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
|
||||
loop(10000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
|
||||
loop(10000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
|
||||
loop(10); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
|
||||
loop(100); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
|
||||
loop(100); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
|
||||
loop(1000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
|
||||
loop(1000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
|
||||
loop(10000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
|
||||
loop(10000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(100); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(100); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(100); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(100); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10);streamsync; streamsync; streamsync; endloop(1);
|
||||
loop(10);streamsync; streamsync; streamsync; endloop(1);
|
||||
loop(100);streamsync; streamsync; streamsync; endloop(1);
|
||||
loop(100);streamsync; streamsync; streamsync; endloop(1);
|
||||
loop(1000);streamsync; streamsync; streamsync; endloop(1);
|
||||
loop(1000);streamsync; streamsync; streamsync; endloop(1);
|
||||
loop(10000);streamsync; streamsync; streamsync; endloop(1);
|
||||
loop(10000);streamsync; streamsync; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(1000); D2H;streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; D2H; D2H; streamsync; endloop(1);
|
||||
loop(10); D2H; D2H; D2H; streamsync; endloop(1);
|
||||
loop(100); D2H; D2H; D2H; streamsync; endloop(1);
|
||||
loop(100); D2H; D2H; D2H; streamsync; endloop(1);
|
||||
loop(1000); D2H; D2H; D2H; streamsync; endloop(1);
|
||||
loop(1000); D2H; D2H; D2H; streamsync; endloop(1);
|
||||
loop(10000); D2H; D2H; D2H; streamsync; endloop(1);
|
||||
loop(10000); D2H; D2H; D2H;streamsync; endloop(1);
|
||||
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; streamsync;H2D;streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10); H2D; streamsync;H2D;streamsync; H2D; streamsync; endloop(1);
|
||||
loop(100); H2D; streamsync;H2D; streamsync;H2D; streamsync; endloop(1);
|
||||
loop(100); H2D;streamsync; H2D; streamsync;H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D;streamsync; H2D;streamsync; H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D;streamsync; H2D; streamsync;H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D;streamsync; H2D; streamsync;H2D; streamsync; endloop(1);
|
||||
loop(10000); H2D;streamsync; H2D; streamsync;H2D; streamsync; endloop(1);
|
||||
loop(10000); H2D;streamsync; H2D;streamsync; H2D; streamsync; endloop(1);
|
||||
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; H2D; H2D; streamsync; endloop(1);
|
||||
loop(10); H2D; H2D; H2D; streamsync; endloop(1);
|
||||
loop(100); H2D; H2D; H2D; streamsync; endloop(1);
|
||||
loop(100); H2D; H2D; H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D; H2D; H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D; H2D; H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D; H2D; H2D; streamsync; endloop(1);
|
||||
loop(10000); H2D; H2D; H2D; streamsync; endloop(1);
|
||||
loop(10000); H2D; H2D; H2D; streamsync; endloop(1);
|
||||
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(10); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(100); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(100); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1);
|
||||
-10
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(10); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(100); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(100); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(10); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(100); NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
|
||||
loop(100); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; NullKernel; NullKernel; NullKernel;NullKernel; streamsync; endloop(1);
|
||||
loop(10); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(100); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
|
||||
loop(100); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1);
|
||||
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); NullKernel; NullKernel; NullKernel; NullKernel;NullKernel;NullKernel; streamsync; endloop(1);
|
||||
loop(10); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
|
||||
loop(100); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;NullKernel;streamsync; endloop(1);
|
||||
loop(100); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
|
||||
loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
|
||||
loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
|
||||
loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; streamsync; endloop(1);
|
||||
loop(10); D2H; streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; H2D; streamsync; endloop(1);
|
||||
loop(10); D2H; H2D; streamsync; endloop(1);
|
||||
loop(100); D2H; H2D; streamsync; endloop(1);
|
||||
loop(100); D2H; H2D; streamsync; endloop(1);
|
||||
loop(1000); D2H; H2D; streamsync; endloop(1);
|
||||
loop(1000); D2H; H2D; streamsync; endloop(1);
|
||||
loop(10000); D2H; H2D; streamsync; endloop(1);
|
||||
loop(10000); D2H; H2D; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(10); D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(100); D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(100); D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); D2H; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); D2H; NullKernel; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10); D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync; H2D; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); D2H; streamsync;NullKernel; streamsync; endloop(1);
|
||||
loop(10); D2H; streamsync;NullKernel; streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync;NullKernel; streamsync; endloop(1);
|
||||
loop(100); D2H; streamsync;NullKernel; streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync;NullKernel; streamsync; endloop(1);
|
||||
loop(1000); D2H; streamsync;NullKernel; streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync;NullKernel; streamsync; endloop(1);
|
||||
loop(10000); D2H; streamsync;NullKernel; streamsync; endloop(1);
|
||||
@@ -1,10 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; streamsync; endloop(1);
|
||||
loop(10); H2D; streamsync; endloop(1);
|
||||
loop(100); H2D; streamsync; endloop(1);
|
||||
loop(100); H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D; streamsync; endloop(1);
|
||||
loop(1000); H2D; streamsync; endloop(1);
|
||||
loop(10000); H2D; streamsync; endloop(1);
|
||||
loop(10000); H2D; streamsync; endloop(1);
|
||||
@@ -1,2 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; D2H; streamsync; endloop(1);
|
||||
loop(10); H2D; D2H; streamsync; endloop(1);
|
||||
loop(100); H2D; D2H; streamsync; endloop(1);
|
||||
loop(100); H2D; D2H; streamsync; endloop(1);
|
||||
loop(1000); H2D; D2H; streamsync; endloop(1);
|
||||
loop(1000); H2D; D2H; streamsync; endloop(1);
|
||||
loop(10000); H2D; D2H; streamsync; endloop(1);
|
||||
loop(10000); H2D; D2H; streamsync; endloop(1);
|
||||
@@ -1,9 +0,0 @@
|
||||
setstream(1);
|
||||
loop(10); H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(10); H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(100); H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(100); H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(1000); H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); H2D; NullKernel; streamsync; endloop(1);
|
||||
loop(10000); H2D; NullKernel; streamsync; endloop(1);
|
||||
Některé soubory nejsou zobrazny, neboť je v této revizi změněno mnoho souborů Zobrazit více
Odkázat v novém úkolu
Zablokovat Uživatele