From 2436c2dd6fba2a4bd188c4282a8d764c7510d2dc Mon Sep 17 00:00:00 2001 From: Sarbojit2019 <52527887+SarbojitAMD@users.noreply.github.com> Date: Fri, 17 Apr 2020 10:31:47 +0530 Subject: [PATCH] [HIPTEST]common changes for unittest (#2017) [ROCm/clr commit: 7808be893f46e450d80085893b9652e1d6424a5d] --- projects/clr/hipamd/CMakeLists.txt | 6 +- .../clr/hipamd/tests/unit/test_common.cpp | 180 +++++++ projects/clr/hipamd/tests/unit/test_common.h | 474 ++++++++++++++++++ 3 files changed, 659 insertions(+), 1 deletion(-) create mode 100644 projects/clr/hipamd/tests/unit/test_common.cpp create mode 100644 projects/clr/hipamd/tests/unit/test_common.h diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index 4061b163ff..b56c47af30 100644 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/projects/clr/hipamd/CMakeLists.txt @@ -552,8 +552,12 @@ if(${RUN_HIT} EQUAL 0) include_directories(${HIP_SRC_PATH}/tests/src) hit_add_directory_recursive(${HIP_SRC_PATH}/tests/src "directed_tests") + # Add unit tests + include_directories(${HIP_SRC_PATH}/tests/unit) + hit_add_directory_recursive(${HIP_SRC_PATH}/tests/unit "unit_tests") + # Add top-level tests to build_tests - add_custom_target(build_tests DEPENDS directed_tests) + add_custom_target(build_tests DEPENDS directed_tests unit_tests) # Add custom target: check add_custom_target(check COMMAND "${CMAKE_COMMAND}" --build . --target test DEPENDS build_tests) diff --git a/projects/clr/hipamd/tests/unit/test_common.cpp b/projects/clr/hipamd/tests/unit/test_common.cpp new file mode 100644 index 0000000000..1c0dcc8c34 --- /dev/null +++ b/projects/clr/hipamd/tests/unit/test_common.cpp @@ -0,0 +1,180 @@ +/* +Copyright (c) 2015-2016 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 "test_common.h" + +// standard global variables that can be set on command line +size_t N = 4 * 1024 * 1024; +char memsetval = 0x42; +int memsetD32val = 0xDEADBEEF; +short memsetD16val = 0xDEAD; +char memsetD8val = 0xDE; +int iterations = 1; +unsigned blocksPerCU = 6; // to hide latency +unsigned threadsPerBlock = 256; +int p_gpuDevice = 0; +unsigned p_verbose = 0; +int p_tests = -1; /*which tests to run. Interpretation is left to each test. default:all*/ +#ifdef _WIN64 +const char* HIP_VISIBLE_DEVICES_STR = "HIP_VISIBLE_DEVICES="; +const char* CUDA_VISIBLE_DEVICES_STR = "CUDA_VISIBLE_DEVICES="; +const char* PATH_SEPERATOR_STR = "\\"; +const char* NULL_DEVICE = "NUL:"; +#else +const char* HIP_VISIBLE_DEVICES_STR = "HIP_VISIBLE_DEVICES"; +const char* CUDA_VISIBLE_DEVICES_STR = "CUDA_VISIBLE_DEVICES"; +const char* PATH_SEPERATOR_STR = "/"; +const char* NULL_DEVICE = "/dev/null"; +#endif + +namespace HipTest { + + +double elapsed_time(long long startTimeUs, long long stopTimeUs) { + return ((double)(stopTimeUs - startTimeUs)) / ((double)(1000)); +} + + +int parseSize(const char* str, size_t* output) { + char* next; + *output = strtoull(str, &next, 0); + int l = strlen(str); + if (l) { + char c = str[l - 1]; // last char. + if ((c == 'k') || (c == 'K')) { + *output *= 1024; + } + if ((c == 'm') || (c == 'M')) { + *output *= (1024 * 1024); + } + if ((c == 'g') || (c == 'G')) { + *output *= (1024 * 1024 * 1024); + } + } + return 1; +} + + +int parseUInt(const char* str, unsigned int* output) { + char* next; + *output = strtoul(str, &next, 0); + return !strlen(next); +} + + +int parseInt(const char* str, int* output) { + char* next; + *output = strtol(str, &next, 0); + return !strlen(next); +} + + +int parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg) { + int extraArgs = 1; + for (int i = 1; i < argc; i++) { + const char* arg = argv[i]; + + if (!strcmp(arg, " ")) { + // skip NULL args. + } else if (!strcmp(arg, "--N") || (!strcmp(arg, "-N"))) { + if (++i >= argc || !HipTest::parseSize(argv[i], &N)) { + failed("Bad N size argument"); + } + } else if (!strcmp(arg, "--threadsPerBlock")) { + if (++i >= argc || !HipTest::parseUInt(argv[i], &threadsPerBlock)) { + failed("Bad threadsPerBlock argument"); + } + } else if (!strcmp(arg, "--blocksPerCU")) { + if (++i >= argc || !HipTest::parseUInt(argv[i], &blocksPerCU)) { + failed("Bad blocksPerCU argument"); + } + } else if (!strcmp(arg, "--memsetval")) { + int ex; + if (++i >= argc || !HipTest::parseInt(argv[i], &ex)) { + failed("Bad memsetval argument"); + } + memsetval = ex; + } else if (!strcmp(arg, "--memsetD32val")) { + int ex; + if (++i >= argc || !HipTest::parseInt(argv[i], &ex)) { + failed("Bad memsetD32val argument"); + } + memsetD32val = ex; + } else if (!strcmp(arg, "--memsetD16val")) { + int ex; + if (++i >= argc || !HipTest::parseInt(argv[i], &ex)) { + failed("Bad memsetD16val argument"); + } + memsetD16val = ex; + } else if (!strcmp(arg, "--memsetD8val")) { + int ex; + if (++i >= argc || !HipTest::parseInt(argv[i], &ex)) { + failed("Bad memsetD8val argument"); + } + memsetD8val = ex; + } else if (!strcmp(arg, "--iterations") || (!strcmp(arg, "-i"))) { + if (++i >= argc || !HipTest::parseInt(argv[i], &iterations)) { + failed("Bad iterations argument"); + } + + } else if (!strcmp(arg, "--gpu") || (!strcmp(arg, "-gpuDevice")) || (!strcmp(arg, "-g"))) { + if (++i >= argc || !HipTest::parseInt(argv[i], &p_gpuDevice)) { + failed("Bad gpuDevice argument"); + } + + } else if (!strcmp(arg, "--verbose") || (!strcmp(arg, "-v"))) { + if (++i >= argc || !HipTest::parseUInt(argv[i], &p_verbose)) { + failed("Bad verbose argument"); + } + } else if (!strcmp(arg, "--tests") || (!strcmp(arg, "-t"))) { + if (++i >= argc || !HipTest::parseInt(argv[i], &p_tests)) { + failed("Bad tests argument"); + } + + } else { + if (failOnUndefinedArg) { + failed("Bad argument '%s'", arg); + } else { + argv[extraArgs++] = argv[i]; + } + } + }; + + return extraArgs; +} + + +unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N) { + int device; + HIPCHECK(hipGetDevice(&device)); + hipDeviceProp_t props; + HIPCHECK(hipGetDeviceProperties(&props, device)); + + unsigned blocks = props.multiProcessorCount * blocksPerCU; + if (blocks * threadsPerBlock > N) { + blocks = (N + threadsPerBlock - 1) / threadsPerBlock; + } + + return blocks; +} + + +} // namespace HipTest diff --git a/projects/clr/hipamd/tests/unit/test_common.h b/projects/clr/hipamd/tests/unit/test_common.h new file mode 100644 index 0000000000..4b55c70164 --- /dev/null +++ b/projects/clr/hipamd/tests/unit/test_common.h @@ -0,0 +1,474 @@ +/* +Copyright (c) 2020-Present 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. +*/ + +/* + * File is intended to C and CPP compliant hence any CPP specic changes + * should be added into CPP section + * + */ + +#ifdef __cplusplus + #include + #include + #if __CUDACC__ + #include + #else + #include + #endif +#endif + +// ************************ GCC section ************************** +#include + +#include "hip/hip_runtime.h" +#include "hip/hip_runtime_api.h" + +#define HC __attribute__((hc)) + + +#define KNRM "\x1B[0m" +#define KRED "\x1B[31m" +#define KGRN "\x1B[32m" +#define KYEL "\x1B[33m" +#define KBLU "\x1B[34m" +#define KMAG "\x1B[35m" +#define KCYN "\x1B[36m" +#define KWHT "\x1B[37m" + +#define passed() \ + printf("%sPASSED!%s\n", KGRN, KNRM); \ + +#define failed(...) \ + printf("%serror: ", KRED); \ + printf(__VA_ARGS__); \ + printf("%s\n",KNRM); \ + return false; + +#define warn(...) \ + printf("%swarn: ", KYEL); \ + printf(__VA_ARGS__); \ + printf("\n"); \ + printf("warn: TEST WARNING\n%s", KNRM); + +#define skipped() printf("%sSkipped subtest %s%s\n",KYEL,__FUNCTION__,KNRM); + +#define HIPCHECK(error) \ + { \ + hipError_t localError = error; \ + if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \ + failed("%serror: '%s'(%d) from %s at %s:%d%s\n", KRED, hipGetErrorString(localError), \ + localError, #error, __FUNCTION__, __LINE__, KNRM); \ + } \ + } + +#define HIPASSERT(condition) \ + if (!(condition)) { \ + failed("%sassertion %s at %s:%d%s \n", KRED, #condition, __FUNCTION__, __LINE__, KNRM); \ + } + + +#define HIPCHECK_API(API_CALL, EXPECTED_ERROR) \ + { \ + hipError_t _e = (API_CALL); \ + if (_e != (EXPECTED_ERROR)) { \ + failed("%sAPI '%s' returned %d(%s) but test expected %d(%s) at %s:%d%s \n", KRED, \ + #API_CALL, _e, hipGetErrorName(_e), EXPECTED_ERROR, \ + hipGetErrorName(EXPECTED_ERROR), __FILE__, __LINE__, KNRM); \ + } \ + } + +#ifdef _WIN64 +#include +#define aligned_alloc(x,y) _aligned_malloc(y,x) +#define aligned_free(x) _aligned_free(x) +#define popen(x,y) _popen(x,y) +#define pclose(x) _pclose(x) +#define setenv(x,y,z) _putenv_s(x,y) +#define unsetenv _putenv +#define fileno(x) _fileno(x) +#define dup(x) _dup(x) +#define dup2(x,y) _dup2(x,y) +#define close(x) _close(x) +#else +#define aligned_free(x) free(x) +#endif + +// standard command-line variables: +extern size_t N; +extern char memsetval; +extern int memsetD32val; +extern short memsetD16val; +extern char memsetD8val; +extern int iterations; +extern unsigned blocksPerCU; +extern unsigned threadsPerBlock; +extern int p_gpuDevice; +extern unsigned p_verbose; +extern int p_tests; +extern const char* HIP_VISIBLE_DEVICES_STR; +extern const char* CUDA_VISIBLE_DEVICES_STR; +extern const char* PATH_SEPERATOR_STR; +extern const char* NULL_DEVICE; + +// ********************* CPP section ********************* +#ifdef __cplusplus + +#ifdef __HIP_PLATFORM_HCC +#define TYPENAME(T) typeid(T).name() +#else +#define TYPENAME(T) "?" +#endif + +namespace HipTest { + +// Returns the current system time in microseconds +inline long long get_time() { +#if __CUDACC__ + struct timeval tv; + gettimeofday(&tv, 0); + return (tv.tv_sec * 1000000) + tv.tv_usec; +#else + return std::chrono::high_resolution_clock::now().time_since_epoch() + /std::chrono::microseconds(1); +#endif +} + +double elapsed_time(long long startTimeUs, long long stopTimeUs); + +int parseSize(const char* str, size_t* output); +int parseUInt(const char* str, unsigned int* output); +int parseInt(const char* str, int* output); +int parseStandardArguments(int argc, char* argv[], bool failOnUndefinedArg); + +unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N); + +template // pointer type +bool checkArray(T hData, T hOutputData, size_t width, size_t height,size_t depth) +{ + for (int i = 0; i < depth; i++) { + for (int j = 0; j < height; j++) { + for (int k = 0; k < width; k++) { + int offset = i*width*height + j*width + k; + if (hData[offset] != hOutputData[offset]) { + std::cerr << '[' << i << ',' << j << ',' << k << "]:" << hData[offset] << "----" << hOutputData[offset]<<" "; + failed("mistmatch at:%d %d %d",i,j,k); + } + } + } + } + return true; +} + +template +bool checkArray(T input, T output, size_t height, size_t width) +{ + for(int i=0; i +__global__ void vectorADD(const T* A_d, const T* B_d, T* C_d, size_t NELEM) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i = offset; i < NELEM; i += stride) { + C_d[i] = A_d[i] + B_d[i]; + } +} + + +template +__global__ void vectorADDReverse(const T* A_d, const T* B_d, T* C_d, + size_t NELEM) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (int64_t i = NELEM - stride + offset; i >= 0; i -= stride) { + C_d[i] = A_d[i] + B_d[i]; + } +} + + +template +__global__ void addCount(const T* A_d, T* C_d, size_t NELEM, int count) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + // Deliberately do this in an inefficient way to increase kernel runtime + for (int i = 0; i < count; i++) { + for (size_t i = offset; i < NELEM; i += stride) { + C_d[i] = A_d[i] + (T)count; + } + } +} + + +template +__global__ void addCountReverse(const T* A_d, T* C_d, int64_t NELEM, int count) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + // Deliberately do this in an inefficient way to increase kernel runtime + for (int i = 0; i < count; i++) { + for (int64_t i = NELEM - stride + offset; i >= 0; i -= stride) { + C_d[i] = A_d[i] + (T)count; + } + } +} + + +template +__global__ void memsetReverse(T* C_d, T val, int64_t NELEM) { + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (int64_t i = NELEM - stride + offset; i >= 0; i -= stride) { + C_d[i] = val; + } +} + + +template +void setDefaultData(size_t numElements, T* A_h, T* B_h, T* C_h) { + // Initialize the host data: + for (size_t i = 0; i < numElements; i++) { + if (A_h) (A_h)[i] = 3.146f + i; // Pi + if (B_h) (B_h)[i] = 1.618f + i; // Phi + if (C_h) (C_h)[i] = 0.0f + i; + } +} + + +template +bool initArraysForHost(T** A_h, T** B_h, T** C_h, size_t N, bool usePinnedHost = false) { + size_t Nbytes = N * sizeof(T); + + if (usePinnedHost) { + if (A_h) { + HIPCHECK(hipHostMalloc((void**)A_h, Nbytes)); + } + if (B_h) { + HIPCHECK(hipHostMalloc((void**)B_h, Nbytes)); + } + if (C_h) { + HIPCHECK(hipHostMalloc((void**)C_h, Nbytes)); + } + } else { + if (A_h) { + *A_h = (T*)malloc(Nbytes); + HIPASSERT(*A_h != NULL); + } + + if (B_h) { + *B_h = (T*)malloc(Nbytes); + HIPASSERT(*B_h != NULL); + } + + if (C_h) { + *C_h = (T*)malloc(Nbytes); + HIPASSERT(*C_h != NULL); + } + } + + setDefaultData(N, A_h ? *A_h : NULL, B_h ? *B_h : NULL, C_h ? *C_h : NULL); + return true; +} + + +template +bool initArrays(T** A_d, T** B_d, T** C_d, T** A_h, T** B_h, T** C_h, size_t N, + bool usePinnedHost = false) { + size_t Nbytes = N * sizeof(T); + + if (A_d) { + HIPCHECK(hipMalloc(A_d, Nbytes)); + } + if (B_d) { + HIPCHECK(hipMalloc(B_d, Nbytes)); + } + if (C_d) { + HIPCHECK(hipMalloc(C_d, Nbytes)); + } + + return initArraysForHost(A_h, B_h, C_h, N, usePinnedHost); +} + + +template +bool freeArraysForHost(T* A_h, T* B_h, T* C_h, bool usePinnedHost) { + if (usePinnedHost) { + if (A_h) { + HIPCHECK(hipHostFree(A_h)); + } + if (B_h) { + HIPCHECK(hipHostFree(B_h)); + } + if (C_h) { + HIPCHECK(hipHostFree(C_h)); + } + } else { + if (A_h) { + free(A_h); + } + if (B_h) { + free(B_h); + } + if (C_h) { + free(C_h); + } + } + return true; +} + +template +bool freeArrays(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHost) { + if (A_d) { + HIPCHECK(hipFree(A_d)); + } + if (B_d) { + HIPCHECK(hipFree(B_d)); + } + if (C_d) { + HIPCHECK(hipFree(C_d)); + } + + return freeArraysForHost(A_h, B_h, C_h, usePinnedHost); +} + +#if defined(__HIP_PLATFORM_HCC__) +template +bool initArrays2DPitch(T** A_d, T** B_d, T** C_d, size_t* pitch_A, size_t* pitch_B, size_t* pitch_C, + size_t numW, size_t numH) { + if (A_d) { + HIPCHECK(hipMallocPitch((void**)A_d, pitch_A, numW * sizeof(T), numH)); + } + if (B_d) { + HIPCHECK(hipMallocPitch((void**)B_d, pitch_B, numW * sizeof(T), numH)); + } + if (C_d) { + HIPCHECK(hipMallocPitch((void**)C_d, pitch_C, numW * sizeof(T), numH)); + } + + HIPASSERT(*pitch_A == *pitch_B); + HIPASSERT(*pitch_A == *pitch_C) + return true; +} + +inline bool initHIPArrays(hipArray** A_d, hipArray** B_d, hipArray** C_d, + const hipChannelFormatDesc* desc, const size_t numW, const size_t numH, + const unsigned int flags) { + if (A_d) { + HIPCHECK(hipMallocArray(A_d, desc, numW, numH, flags)); + } + if (B_d) { + HIPCHECK(hipMallocArray(B_d, desc, numW, numH, flags)); + } + if (C_d) { + HIPCHECK(hipMallocArray(C_d, desc, numW, numH, flags)); + } + return true; +} +#endif + +// Assumes C_h contains vector add of A_h + B_h +// Calls the test "failed" macro if a mismatch is detected. +template +size_t checkVectorADD(T* A_h, T* B_h, T* result_H, size_t N, bool expectMatch = true, + bool reportMismatch = true) { + size_t mismatchCount = 0; + size_t firstMismatch = 0; + size_t mismatchesToPrint = 10; + for (size_t i = 0; i < N; i++) { + T expected = A_h[i] + B_h[i]; + if (result_H[i] != expected) { + if (mismatchCount == 0) { + firstMismatch = i; + } + mismatchCount++; + if ((mismatchCount <= mismatchesToPrint) && expectMatch) { + std::cout << std::fixed << std::setprecision(32); + std::cout << "At " << i << std::endl; + std::cout << " Computed:" << result_H[i] << std::endl; + std::cout << " Expected:" << expected << std::endl; + } + } + } + + if (reportMismatch) { + if (expectMatch) { + if (mismatchCount) { + failed("%zu mismatches ; first at index:%zu\n", mismatchCount, firstMismatch); + } + } else { + if (mismatchCount == 0) { + failed("expected mismatches but did not detect any!"); + } + } + } + + return mismatchCount; +} + + +// Assumes C_h contains vector add of A_h + B_h +// Calls the test "failed" macro if a mismatch is detected. +template +bool checkTest(T* expected_H, T* result_H, size_t N, bool expectMatch = true) { + size_t mismatchCount = 0; + size_t firstMismatch = 0; + size_t mismatchesToPrint = 10; + for (size_t i = 0; i < N; i++) { + if (result_H[i] != expected_H[i]) { + if (mismatchCount == 0) { + firstMismatch = i; + } + mismatchCount++; + if ((mismatchCount <= mismatchesToPrint) && expectMatch) { + std::cout << std::fixed << std::setprecision(32); + std::cout << "At " << i << std::endl; + std::cout << " Computed:" << result_H[i] << std::endl; + std::cout << " Expected:" << expected_H[i] << std::endl; + } + } + } + + if (expectMatch) { + if (mismatchCount) { + fprintf(stderr, "%zu mismatches ; first at index:%zu\n", mismatchCount, firstMismatch); + // failed("%zu mismatches ; first at index:%zu\n", mismatchCount, firstMismatch); + } + } else { + if (mismatchCount == 0) { + failed("expected mismatches but did not detect any!"); + } + } + return true; +} + +}; // namespace HipTest +#endif //__cplusplus