AI/rocm-systems
2
0
انشعاب 0
کد مسائل تقاضاهای واکشی Actions Packages پروژه‌ها انتشارها دانشنامه فعالیت

SWDEV-292393 - [catch2][dtest] Tests for hipMemcpy related apis.

فهرست منبع 
Migrated all hipMemcpy related APIs to CATCH2 framework by optmizing
the code and moving the stress related tests to stress folder.

Change-Id: Id47669b49304c35d1a68fabdaaf3f6e3ab0428a5


[ROCm/hip commit: 346a77b4c0]
This commit is contained in:
DURGESH KROTTAPALLI
2021-07-06 23:48:24 +05:30
والد e9c8e3dd04
کامیت 7e9e6f1729
20فایلهای تغییر یافته به همراه3593 افزوده شده و 62 حذف شده
دانلود پرونده وصله دانلود فایل تغییرات diff Expand all files Collapse all files
projects/hip/tests/catch/include/hip_test_checkers.hh
+41 -37
مشاهده پرونده
@@ -1,14 +1,13 @@
#pragma once
#include "hip_test_common.hh"
#include <iostream>
using namespace std;
#define guarantee(cond, str) \
{ \
if (!(cond)) { \
std::cout << str << std::endl; \
abort(); \
} \
}
#define guarantee(cond, str) \
{ \
if (!(cond)) { \
INFO("guarantee failed: " << str); \
abort(); \
} \
}
namespace HipTest {
@@ -49,31 +48,20 @@ size_t checkVectors(T* A, T* B, T* Out, size_t N, T (*F)(T a, T b), bool expectM
return mismatchCount;
}
template<typename T> // pointer type
void checkArray(T hData, T hOutputData, size_t width, size_t height,size_t depth) {
bool checkArray(T* hData, T* hOutputData, size_t width, size_t height,size_t depth = 1) {
for (size_t i = 0; i < depth; i++) {
for (size_t j = 0; j < height; j++) {
for (size_t k = 0; k < width; k++) {
int offset = i*width*height + j*width + k;
if (hData[offset] != hOutputData[offset]) {
cerr << '[' << i << ',' << j << ',' << k << "]:" << hData[offset] << "----"
<< hOutputData[offset]<<" ";
cout << "mistmatch at: " << i<< j<<k;
INFO("Mismatch at [" << i << "," << j << "," << k << "]:"
<< hData[offset] << "----" << hOutputData[offset]);
CHECK(false);
return false;
}
}
}
}
}
template<typename T> // pointer type
bool checkArray(T *result, T *compare, size_t width, size_t height) {
for (size_t i = 0; i < height; i++) {
for (size_t j = 0; j < width; j++) {
if (result[(i*width) + j] != compare[(i*width) + j]) {
std::cout << result[(i*width) + j] << "\t" << compare[(i*width) + j] << std::endl;
return false;
}
}
}
return true;
}
@@ -103,17 +91,17 @@ template <typename T> void setDefaultData(size_t numElements, T* A_h, T* B_h, T*
for (size_t i = 0; i < numElements; i++) {
if (std::is_same<T, int>::value || std::is_same<T, unsigned int>::value) {
if (A_h) (A_h)[i] = 3;
if (B_h) (B_h)[i] = 4;
if (C_h) (C_h)[i] = 5;
if (A_h) A_h[i] = 3;
if (B_h) B_h[i] = 4;
if (C_h) C_h[i] = 5;
} else if(std::is_same<T, char>::value || std::is_same<T, unsigned char>::value) {
if (A_h) (A_h)[i] = 'a';
if (B_h) (B_h)[i] = 'b';
if (C_h) (C_h)[i] = 'c';
if (A_h) A_h[i] = 'a';
if (B_h) B_h[i] = 'b';
if (C_h) C_h[i] = 'c';
} else {
if (A_h) (A_h)[i] = 3.146f + i;
if (B_h) (B_h)[i] = 1.618f + i;
if (C_h) (C_h)[i] = 1.4f + i;
if (A_h) A_h[i] = 3.146f + i;
if (B_h) B_h[i] = 1.618f + i;
if (C_h) C_h[i] = 1.4f + i;
}
}
}
@@ -135,21 +123,21 @@ bool initArraysForHost(T** A_h, T** B_h, T** C_h, size_t N, bool usePinnedHost =
} else {
if (A_h) {
*A_h = (T*)malloc(Nbytes);
REQUIRE(*A_h != NULL);
REQUIRE(*A_h != nullptr);
}
if (B_h) {
*B_h = (T*)malloc(Nbytes);
REQUIRE(*B_h != NULL);
REQUIRE(*B_h != nullptr);
}
if (C_h) {
*C_h = (T*)malloc(Nbytes);
REQUIRE(*C_h != NULL);
REQUIRE(*C_h != nullptr);
}
}
setDefaultData(N, A_h ? *A_h : NULL, B_h ? *B_h : NULL, C_h ? *C_h : NULL);
setDefaultData(N, A_h ? *A_h : nullptr, B_h ? *B_h : nullptr, C_h ? *C_h : nullptr);
return true;
}
@@ -210,4 +198,20 @@ bool freeArrays(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHo
return freeArraysForHost(A_h, B_h, C_h, usePinnedHost);
}
template <typename T>
unsigned setNumBlocks(T blocksPerCU, T threadsPerBlock,
size_t N) {
int device;
HIP_CHECK(hipGetDevice(&device));
hipDeviceProp_t props;
HIP_CHECK(hipGetDeviceProperties(&props, device));
unsigned blocks = props.multiProcessorCount * blocksPerCU;
if (blocks * threadsPerBlock > N) {
blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
}
return blocks;
}
} // namespace HipTest
projects/hip/tests/catch/include/hip_test_common.hh
+18
مشاهده پرونده
@@ -55,6 +55,24 @@ THE SOFTWARE.
#include <chrono>
#endif
#define HIPCHECK(error) \
{ \
hipError_t localError = error; \
if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \
printf("error: '%s'(%d) from %s at %s:%d\n", hipGetErrorString(localError), \
localError, #error, __FILE__, __LINE__); \
abort(); \
} \
}
#define HIPASSERT(condition) \
if (!(condition)) { \
printf("assertion %s at %s:%d \n", #condition, __FILE__, __LINE__); \
abort(); \
}
// Utility Functions
namespace HipTest {
static inline int getDeviceCount() {
projects/hip/tests/catch/include/hip_test_kernels.hh
+9 -1
مشاهده پرونده
@@ -72,7 +72,6 @@ __global__ void addCountReverse(const T* A_d, T* C_d, int64_t NELEM, int count)
}
}
template <typename T> __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;
@@ -81,4 +80,13 @@ template <typename T> __global__ void memsetReverse(T* C_d, T val, int64_t NELEM
C_d[i] = val;
}
}
template <typename T> __global__ void vector_square(const T* A_d, T* C_d, size_t N_ELMTS) {
size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
for (size_t i = gputhread; i < N_ELMTS; i += stride) {
C_d[i] = A_d[i] * A_d[i];
}
}
} // namespace HipTest