AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

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

Browse Source 
Migrated all hipMemcpy related APIs to CATCH2 framework by optmizing
the code and moving the stress related tests to stress folder.

Change-Id: Id47669b49304c35d1a68fabdaaf3f6e3ab0428a5
This commit is contained in:
DURGESH KROTTAPALLI
2021-07-06 23:48:24 +05:30
parent 8093223eec
commit 18591bc68f
20 changed files with 3593 additions and 62 deletions
Download Patch File Download Diff File Expand all files Collapse all files
catch/hipTestMain/CMakeLists.txt
+15
View File
@@ -54,3 +54,18 @@ if(UNIX)
catch_discover_tests(MultiProcTests PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST")
add_dependencies(build_tests MultiProcTests)
endif()
add_executable(StressTest EXCLUDE_FROM_ALL main.cc hip_test_context.cc)
add_custom_target(build_stress_test)
if(HIP_PLATFORM MATCHES "amd")
set_property(TARGET StressTest PROPERTY CXX_STANDARD 17)
else()
target_compile_options(StressTest PUBLIC -std=c++17)
endif()
if(HIP_PLATFORM MATCHES "amd")
target_link_libraries(StressTest PRIVATE printf stream)
endif()
target_link_libraries(StressTest PRIVATE memory stdc++fs)
add_dependencies(build_stress_test StressTest)
add_custom_target(stress_test COMMAND StressTest)
catch/include/hip_test_checkers.hh
+41 -37
View File
@@ -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
catch/include/hip_test_common.hh
+18
View File
@@ -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() {
catch/include/hip_test_kernels.hh
+9 -1
View File
@@ -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
catch/multiproc/hipMallocConcurrency.cc
+4 -4
View File
@@ -12,10 +12,10 @@
#include <hip_test_common.hh>
size_t N = 4 * 1024 * 1024;
unsigned blocksPerCU = 6; // to hide latency
unsigned threadsPerBlock = 256;
static constexpr size_t N = 4 * 1024 * 1024;
static constexpr unsigned blocksPerCU = 6; // to hide latency
static constexpr unsigned threadsPerBlock = 256;
/**
* Validates data consitency on supplied gpu
*/
@@ -31,7 +31,7 @@ bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
printf("tgs allocating..\n");
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
catch/stress/CMakeLists.txt
-8
View File
@@ -1,13 +1,5 @@
add_custom_target(build_stress_test)
add_executable(StressTest EXCLUDE_FROM_ALL ../hipTestMain/main.cc ../hipTestMain/hip_test_context.cc)
set_property(TARGET StressTest PROPERTY CXX_STANDARD 17)
target_link_libraries(StressTest PRIVATE stdc++fs)
add_dependencies(build_stress_test StressTest)
add_custom_target(stress_test COMMAND StressTest)
add_subdirectory(memory)
if(HIP_PLATFORM MATCHES "amd")
add_subdirectory(printf)
add_subdirectory(stream)
target_link_libraries(StressTest PRIVATE printf stream)
endif()
target_link_libraries(StressTest PRIVATE memory)
catch/stress/memory/CMakeLists.txt
+1
View File
@@ -1,6 +1,7 @@
# Common Tests - Test independent of all platforms
set(TEST_SRC
memcpy.cc
hipMemcpyMThreadMSize.cc
)
# Create shared lib of all tests
catch/stress/memory/hipMemcpyMThreadMSize.cc
+275
View File
@@ -0,0 +1,275 @@
/*
Copyright (c) 2021 - 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.
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
#include <utility>
#include <vector>
/*
This testfile verifies the following scenarios of all hipMemcpy API
1. Multi thread
2. Multi size
*/
static auto Available_Gpus{0};
static constexpr auto MAX_GPU{256};
enum apiToTest {TEST_MEMCPY, TEST_MEMCPYH2D, TEST_MEMCPYD2H, TEST_MEMCPYD2D,
TEST_MEMCPYASYNC, TEST_MEMCPYH2DASYNC, TEST_MEMCPYD2HASYNC,
TEST_MEMCPYD2DASYNC};
template<typename TestType>
void Memcpy_And_verify(int NUM_ELM) {
TestType *A_h, *B_h;
for (apiToTest api = TEST_MEMCPY; api <= TEST_MEMCPYD2DASYNC;
api = apiToTest(api + 1)) {
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&A_h, &B_h, nullptr,
NUM_ELM);
HIP_CHECK(hipGetDeviceCount(&Available_Gpus));
TestType *A_d[MAX_GPU];
hipStream_t stream[MAX_GPU];
for (int i = 0; i < Available_Gpus; ++i) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipMalloc(&A_d[i], NUM_ELM * sizeof(TestType)));
if (api >= TEST_MEMCPYD2D) {
HIP_CHECK(hipStreamCreate(&stream[i]));
}
}
HIP_CHECK(hipSetDevice(0));
int canAccessPeer = 0;
switch (api) {
case TEST_MEMCPY:
{
// To test hipMemcpy()
// Copying data from host to individual devices followed by copying
// back to host and verifying the data consistency.
for (int i = 0; i < Available_Gpus; ++i) {
HIP_CHECK(hipMemcpy(A_d[i], A_h, NUM_ELM * sizeof(TestType),
hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(TestType),
hipMemcpyDeviceToHost));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
// Device to Device copying for all combinations
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
if (canAccessPeer) {
HIP_CHECK(hipMemcpy(A_d[j], A_d[i], NUM_ELM * sizeof(TestType),
hipMemcpyDefault));
// Copying in reverse dir of above to check if bidirectional
// access is happening without any error
HIP_CHECK(hipMemcpy(A_d[i], A_d[j], NUM_ELM * sizeof(TestType),
hipMemcpyDefault));
// Copying data to host to verify the content
HIP_CHECK(hipMemcpy(B_h, A_d[j], NUM_ELM * sizeof(TestType),
hipMemcpyDefault));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
}
}
break;
}
case TEST_MEMCPYH2D: // To test hipMemcpyHtoD()
{
for (int i = 0; i < Available_Gpus; ++i) {
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[i]),
A_h, NUM_ELM * sizeof(TestType)));
// Copying data from device to host to check data consistency
HIP_CHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(TestType),
hipMemcpyDeviceToHost));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
break;
}
case TEST_MEMCPYD2H: // To test hipMemcpyDtoH()--done
{
for (int i = 0; i < Available_Gpus; ++i) {
HIP_CHECK(hipMemcpy(A_d[i], A_h, NUM_ELM * sizeof(TestType),
hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpyDtoH(B_h, hipDeviceptr_t(A_d[i]),
NUM_ELM * sizeof(TestType)));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
break;
}
case TEST_MEMCPYD2D: // To test hipMemcpyDtoD()
{
if (Available_Gpus > 1) {
// First copy data from H to D and then
// from D to D followed by D to H
// HIP_CHECK(hipMemcpyHtoD(A_d[0], A_h,
// NUM_ELM * sizeof(TestType)));
int canAccessPeer = 0;
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
if (canAccessPeer) {
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[i]),
A_h, NUM_ELM * sizeof(TestType)));
HIP_CHECK(hipMemcpyDtoD(hipDeviceptr_t(A_d[j]),
hipDeviceptr_t(A_d[i]), NUM_ELM * sizeof(TestType)));
// Copying in direction reverse of above to check if
// bidirectional
// access is happening without any error
HIP_CHECK(hipMemcpyDtoD(hipDeviceptr_t(A_d[i]),
hipDeviceptr_t(A_d[j]), NUM_ELM * sizeof(TestType)));
HIP_CHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(TestType),
hipMemcpyDeviceToHost));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
}
}
} else {
// As DtoD is not possible transfer data from HtH(A_h to B_h)
// so as to get through verification step
HIP_CHECK(hipMemcpy(B_h, A_h, NUM_ELM * sizeof(TestType),
hipMemcpyHostToHost));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
break;
}
case TEST_MEMCPYASYNC:
{
// To test hipMemcpyAsync()
// Copying data from host to individual devices followed by copying
// back to host and verifying the data consistency.
for (int i = 0; i < Available_Gpus; ++i) {
HIP_CHECK(hipMemcpyAsync(A_d[i], A_h, NUM_ELM * sizeof(TestType),
hipMemcpyHostToDevice, stream[i]));
HIP_CHECK(hipMemcpyAsync(B_h, A_d[i], NUM_ELM * sizeof(TestType),
hipMemcpyDeviceToHost, stream[i]));
HIP_CHECK(hipStreamSynchronize(stream[i]));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
// Device to Device copying for all combinations
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
if (canAccessPeer) {
HIP_CHECK(hipMemcpyAsync(A_d[j], A_d[i],
NUM_ELM * sizeof(TestType),
hipMemcpyDefault, stream[i]));
// Copying in direction reverse of above to
// check if bidirectional
// access is happening without any error
HIP_CHECK(hipMemcpyAsync(A_d[i], A_d[j],
NUM_ELM * sizeof(TestType),
hipMemcpyDefault, stream[i]));
HIP_CHECK(hipStreamSynchronize(stream[i]));
HIP_CHECK(hipMemcpy(B_h, A_d[j], NUM_ELM * sizeof(TestType),
hipMemcpyDefault));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
}
}
break;
}
case TEST_MEMCPYH2DASYNC: // To test hipMemcpyHtoDAsync()
{
for (int i = 0; i < Available_Gpus; ++i) {
HIP_CHECK(hipMemcpyHtoDAsync(hipDeviceptr_t(A_d[i]), A_h,
NUM_ELM * sizeof(TestType), stream[i]));
HIP_CHECK(hipStreamSynchronize(stream[i]));
// Copying data from device to host to check data consistency
HIP_CHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(TestType),
hipMemcpyDeviceToHost));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
break;
}
case TEST_MEMCPYD2HASYNC: // To test hipMemcpyDtoHAsync()
{
for (int i = 0; i < Available_Gpus; ++i) {
HIP_CHECK(hipMemcpy(A_d[i], A_h, NUM_ELM * sizeof(TestType),
hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpyDtoHAsync(B_h, hipDeviceptr_t(A_d[i]),
NUM_ELM * sizeof(TestType), stream[i]));
HIP_CHECK(hipStreamSynchronize(stream[i]));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
break;
}
case TEST_MEMCPYD2DASYNC: // To test hipMemcpyDtoDAsync()
{
if (Available_Gpus > 1) {
// First copy data from H to D and then from D to D followed by D2H
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[0]),
A_h, NUM_ELM * sizeof(TestType)));
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(j));
HIP_CHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d[j]),
hipDeviceptr_t(A_d[i]), NUM_ELM * sizeof(TestType),
stream[i]));
// Copying in direction reverse of above to check if
// bidirectional
// access is happening without any error
HIP_CHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d[i]),
hipDeviceptr_t(A_d[j]), NUM_ELM * sizeof(TestType),
stream[i]));
HIP_CHECK(hipStreamSynchronize(stream[i]));
HIP_CHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(TestType),
hipMemcpyDeviceToHost));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
}
}
} else {
// As DtoD is not possible we will transfer data
// from HtH(A_h to B_h)
// so as to get through verification step
HIP_CHECK(hipMemcpy(B_h, A_h, NUM_ELM * sizeof(TestType),
hipMemcpyHostToHost));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
break;
}
}
for (int i = 0; i < Available_Gpus; ++i) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipFree((A_d[i])));
if (api >= TEST_MEMCPYD2D) {
HIP_CHECK(hipStreamDestroy(stream[i]));
}
}
HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr,
A_h, B_h, nullptr, false);
}
}
TEMPLATE_TEST_CASE("Unit_hipMemcpy_multiDevice-AllAPIs", "",
char, int, size_t, long double) {
auto diff_size = GENERATE(1, 5, 10, 100, 1024, 10*1024, 100*1024,
1024*1024, 10*1024*1024, 100*1024*1024,
1024*1024*1024);
size_t free = 0, total = 0;
HIP_CHECK(hipMemGetInfo(&free, &total));
if ((diff_size * sizeof(TestType)) <= free) {
Memcpy_And_verify<TestType>(diff_size);
HIP_CHECK(hipDeviceSynchronize());
}
}
catch/unit/memory/CMakeLists.txt
+8
View File
@@ -16,6 +16,14 @@ set(TEST_SRC
hipMemcpy2DFromArrayAsync.cc
hipMemcpyAtoH.cc
hipMemcpyHtoA.cc
hipMemcpyDtoD.cc
hipMemcpyDtoDAsync.cc
hipMemcpyAsync.cc
hipMemcpy.cc
hipMemcpyWithStream.cc
hipMemcpyAllApiNegative.cc
hipMemcpyWithStreamMultiThread.cc
hipMemcpy_MultiThread.cc
)
# Create shared lib of all tests
catch/unit/memory/hipMemcpy.cc
+645
View File
@@ -0,0 +1,645 @@
/*
Copyright (c) 2021 - 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.
*/
/*
This testcase verifies following scenarios
1. hipMemcpy API along with kernel launch with different data types
2. H2D-D2D-D2H scenarios for unpinned and pinned memory
3. Boundary checks with different sizes
4. Multithread scenario
5. device offset scenario
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
#ifdef _WIN32
#define WIN32_LEAN_AND_MEAN
#include <windows.h>
#else
#include "sys/types.h"
#include "sys/sysinfo.h"
#endif
static constexpr auto NUM_ELM{4*1024 * 1024};
static unsigned blocksPerCU{6}; // to hide latency
static unsigned threadsPerBlock{256};
template<typename T>
class DeviceMemory {
public:
explicit DeviceMemory(size_t numElements);
DeviceMemory() = delete;
~DeviceMemory();
T* A_d() const { return _A_d + _offset; }
T* B_d() const { return _B_d + _offset; }
T* C_d() const { return _C_d + _offset; }
T* C_dd() const { return _C_dd + _offset; }
size_t maxNumElements() const { return _maxNumElements; }
void offset(int offset) { _offset = offset; }
int offset() const { return _offset; }
private:
T* _A_d;
T* _B_d;
T* _C_d;
T* _C_dd;
size_t _maxNumElements;
int _offset;
};
template <typename T>
DeviceMemory<T>::DeviceMemory(size_t numElements) :
_maxNumElements(numElements), _offset(0) {
T** np = nullptr;
HipTest::initArrays(&_A_d, &_B_d, &_C_d, np, np, np, numElements, 0);
size_t sizeElements = numElements * sizeof(T);
HIP_CHECK(hipMalloc(&_C_dd, sizeElements));
}
template <typename T>
DeviceMemory<T>::~DeviceMemory() {
T* np = nullptr;
HipTest::freeArrays<T>(_A_d, _B_d, _C_d, np, np, np, 0);
HIP_CHECK(hipFree(_C_dd));
_C_dd = NULL;
}
template <typename T>
class HostMemory {
public:
HostMemory(size_t numElements, bool usePinnedHost);
HostMemory() = delete;
void reset(size_t numElements, bool full = false);
~HostMemory();
T* A_h() const { return _A_h + _offset; }
T* B_h() const { return _B_h + _offset; }
T* C_h() const { return _C_h + _offset; }
size_t maxNumElements() const { return _maxNumElements; }
void offset(int offset) { _offset = offset; }
int offset() const { return _offset; }
// Host arrays, secondary copy
T* A_hh;
T* B_hh;
bool _usePinnedHost;
private:
size_t _maxNumElements;
int _offset;
// Host arrays
T* _A_h;
T* _B_h;
T* _C_h;
};
template <typename T>
HostMemory<T>::HostMemory(size_t numElements, bool usePinnedHost)
: _usePinnedHost(usePinnedHost), _maxNumElements(numElements), _offset(0) {
T** np = nullptr;
HipTest::initArrays(np, np, np, &_A_h, &_B_h, &_C_h,
numElements, usePinnedHost);
A_hh = NULL;
B_hh = NULL;
size_t sizeElements = numElements * sizeof(T);
if (usePinnedHost) {
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_hh), sizeElements,
hipHostMallocDefault));
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&B_hh), sizeElements,
hipHostMallocDefault));
} else {
A_hh = reinterpret_cast<T*>(malloc(sizeElements));
B_hh = reinterpret_cast<T*>(malloc(sizeElements));
}
}
template <typename T>
void HostMemory<T>::reset(size_t numElements, bool full) {
// Initialize the host data:
for (size_t i = 0; i < numElements; i++) {
(A_hh)[i] = 1097.0 + i;
(B_hh)[i] = 1492.0 + i; // Phi
if (full) {
(_A_h)[i] = 3.146f + i; // Pi
(_B_h)[i] = 1.618f + i; // Phi
}
}
}
template <typename T>
HostMemory<T>::~HostMemory() {
HipTest::freeArraysForHost(_A_h, _B_h, _C_h, _usePinnedHost);
if (_usePinnedHost) {
HIP_CHECK(hipHostFree(A_hh));
HIP_CHECK(hipHostFree(B_hh));
} else {
free(A_hh);
free(B_hh);
}
}
#ifdef _WIN32
void memcpytest2_get_host_memory(size_t *free, size_t *total) {
MEMORYSTATUSEX status;
status.dwLength = sizeof(status);
GlobalMemoryStatusEx(&status);
// Windows doesn't allow allocating more than half of system memory to the gpu
// Since the runtime also needs space for its internal allocations,
// we should not try to allocate more than 40% of reported system memory,
// otherwise we can run into OOM issues.
*free = static_cast<size_t>(0.4 * status.ullAvailPhys);
*total = static_cast<size_t>(0.4 * status.ullTotalPhys);
}
#else
struct sysinfo memInfo;
void memcpytest2_get_host_memory(size_t *free, size_t *total) {
sysinfo(&memInfo);
uint64_t freePhysMem = memInfo.freeram;
freePhysMem *= memInfo.mem_unit;
*free = freePhysMem;
uint64_t totalPhysMem = memInfo.totalram;
totalPhysMem *= memInfo.mem_unit;
*total = totalPhysMem;
}
#endif
//---
// Test many different kinds of memory copies.
// The subroutine allocates memory , copies to device, runs a vector
// add kernel, copies back, and
// checks the result.
//
// IN: numElements controls the number of elements used for allocations.
// IN: usePinnedHost : If true, allocate host with hipHostMalloc and is pinned
// else allocate host
// memory with malloc. IN: useHostToHost : If true, add an extra
// host-to-host copy. IN:
// useDeviceToDevice : If true, add an extra deviceto-device copy after
// result is produced. IN:
// useMemkindDefault : If true, use memkinddefault
// (runtime figures out direction). if false, use
// explicit memcpy direction.
//
template <typename T>
void memcpytest2(DeviceMemory<T>* dmem, HostMemory<T>* hmem,
size_t numElements, bool useHostToHost,
bool useDeviceToDevice, bool useMemkindDefault) {
size_t sizeElements = numElements * sizeof(T);
hmem->reset(numElements);
assert(numElements <= dmem->maxNumElements());
assert(numElements <= hmem->maxNumElements());
if (useHostToHost) {
// Do some extra host-to-host copies here to mix things up:
HIP_CHECK(hipMemcpy(hmem->A_hh, hmem->A_h(), sizeElements,
useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToHost));
HIP_CHECK(hipMemcpy(hmem->B_hh, hmem->B_h(), sizeElements,
useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToHost));
HIP_CHECK(hipMemcpy(dmem->A_d(), hmem->A_hh, sizeElements,
useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(dmem->B_d(), hmem->B_hh, sizeElements,
useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice));
} else {
HIP_CHECK(hipMemcpy(dmem->A_d(), hmem->A_h(), sizeElements,
useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(dmem->B_d(), hmem->B_h(), sizeElements,
useMemkindDefault ? hipMemcpyDefault : hipMemcpyHostToDevice));
}
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0,
static_cast<const T*>(dmem->A_d()), static_cast<const T*>(dmem->B_d()),
dmem->C_d(), numElements);
if (useDeviceToDevice) {
// Do an extra device-to-device copy here to mix things up:
HIP_CHECK(hipMemcpy(dmem->C_dd(), dmem->C_d(), sizeElements,
useMemkindDefault ? hipMemcpyDefault : hipMemcpyDeviceToDevice));
// Destroy the original dmem->C_d():
HIP_CHECK(hipMemset(dmem->C_d(), 0x5A, sizeElements));
HIP_CHECK(hipMemcpy(hmem->C_h(), dmem->C_dd(), sizeElements,
useMemkindDefault ? hipMemcpyDefault : hipMemcpyDeviceToHost));
} else {
HIP_CHECK(hipMemcpy(hmem->C_h(), dmem->C_d(), sizeElements,
useMemkindDefault ? hipMemcpyDefault : hipMemcpyDeviceToHost));
}
HIP_CHECK(hipDeviceSynchronize());
HipTest::checkVectorADD(hmem->A_h(), hmem->B_h(), hmem->C_h(), numElements);
printf(" %s success\n", __func__);
}
// Try all the 16 possible combinations to memcpytest2 - usePinnedHost,
// useHostToHost,
// useDeviceToDevice, useMemkindDefault
template <typename T>
void memcpytest2_for_type(size_t numElements) {
DeviceMemory<T> memD(numElements);
HostMemory<T> memU(numElements, 0 /*usePinnedHost*/);
HostMemory<T> memP(numElements, 1 /*usePinnedHost*/);
for (int usePinnedHost = 0; usePinnedHost <= 1; usePinnedHost++) {
for (int useHostToHost = 0; useHostToHost <= 1; useHostToHost++) {
for (int useDeviceToDevice = 0; useDeviceToDevice <= 1;
useDeviceToDevice++) {
for (int useMemkindDefault = 0; useMemkindDefault <= 1;
useMemkindDefault++) {
memcpytest2<T>(&memD, usePinnedHost ? &memP : &memU,
numElements, useHostToHost,
useDeviceToDevice, useMemkindDefault);
}
}
}
}
}
// Try many different sizes to memory copy.
template <typename T>
void memcpytest2_sizes(size_t maxElem = 0) {
int deviceId;
HIP_CHECK(hipGetDevice(&deviceId));
size_t free, total, freeCPU, totalCPU;
HIP_CHECK(hipMemGetInfo(&free, &total));
memcpytest2_get_host_memory(&freeCPU, &totalCPU);
if (maxElem == 0) {
// Use lesser maxElem if not enough host memory available
size_t maxElemGPU = free / sizeof(T) / 8;
size_t maxElemCPU = freeCPU / sizeof(T) / 8;
maxElem = maxElemGPU < maxElemCPU ? maxElemGPU : maxElemCPU;
}
HIP_CHECK(hipDeviceReset());
DeviceMemory<T> memD(maxElem);
HostMemory<T> memU(maxElem, 0 /*usePinnedHost*/);
HostMemory<T> memP(maxElem, 1 /*usePinnedHost*/);
for (size_t elem = 1; elem <= maxElem; elem *= 2) {
memcpytest2<T>(&memD, &memU, elem, 1, 1, 0); // unpinned host
memcpytest2<T>(&memD, &memP, elem, 1, 1, 0); // pinned host
}
}
// Try many different sizes to memory copy.
template <typename T>
void memcpytest2_offsets(size_t maxElem, bool devOffsets, bool hostOffsets) {
int deviceId;
HIP_CHECK(hipGetDevice(&deviceId));
size_t free, total;
HIP_CHECK(hipMemGetInfo(&free, &total));
HIP_CHECK(hipDeviceReset());
DeviceMemory<T> memD(maxElem);
HostMemory<T> memU(maxElem, 0 /*usePinnedHost*/);
HostMemory<T> memP(maxElem, 1 /*usePinnedHost*/);
size_t elem = maxElem / 2;
for (size_t offset = 0; offset < 512; offset++) {
assert(elem + offset < maxElem);
if (devOffsets) {
memD.offset(offset);
}
if (hostOffsets) {
memU.offset(offset);
memP.offset(offset);
}
memcpytest2<T>(&memD, &memU, elem, 1, 1, 0); // unpinned host
memcpytest2<T>(&memD, &memP, elem, 1, 1, 0); // pinned host
}
for (size_t offset = 512; offset < elem; offset *= 2) {
assert(elem + offset < maxElem);
if (devOffsets) {
memD.offset(offset);
}
if (hostOffsets) {
memU.offset(offset);
memP.offset(offset);
}
memcpytest2<T>(&memD, &memU, elem, 1, 1, 0); // unpinned host
memcpytest2<T>(&memD, &memP, elem, 1, 1, 0); // pinned host
}
}
// Create multiple threads to stress multi-thread locking behavior in the
// allocation/deallocation/tracking logic:
template <typename T>
void multiThread_1(bool serialize, bool usePinnedHost) {
DeviceMemory<T> memD(NUM_ELM);
HostMemory<T> mem1(NUM_ELM, usePinnedHost);
HostMemory<T> mem2(NUM_ELM, usePinnedHost);
std::thread t1(memcpytest2<T>, &memD, &mem1, NUM_ELM, 0, 0, 0);
if (serialize) {
t1.join();
}
std::thread t2(memcpytest2<T>, &memD, &mem2, NUM_ELM, 0, 0, 0);
if (serialize) {
t2.join();
}
}
/*
This testcase verifies hipMemcpy API
Initializes device variables
Launches kernel and performs the sum of device variables
copies the result to host variable and validates the result.
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpy_KernelLaunch", "", int, float,
double) {
size_t Nbytes = NUM_ELM * sizeof(TestType);
TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, NUM_ELM, false);
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0,
static_cast<const TestType*>(A_d),
static_cast<const TestType*>(B_d), C_d, NUM_ELM);
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipDeviceSynchronize());
HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM);
HipTest::freeArrays<TestType>(A_d, B_d, C_d, A_h, B_h, C_h, false);
}
/*
This testcase verifies the following scenarios
1. H2H,H2PinMem and PinnedMem2Host
2. H2D-D2D-D2H in same GPU
3. Pinned Host Memory to device variables in same GPU
4. Device context change
5. H2D-D2D-D2H peer GPU
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpy_H2H-H2D-D2H-H2PinMem", "", int,
float, double) {
TestType *A_d{nullptr}, *B_d{nullptr};
TestType *A_h{nullptr}, *B_h{nullptr};
TestType *A_Ph{nullptr}, *B_Ph{nullptr};
HIP_CHECK(hipSetDevice(0));
HipTest::initArrays<TestType>(&A_d, &B_d, nullptr,
&A_h, &B_h, nullptr,
NUM_ELM*sizeof(TestType));
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&A_Ph, &B_Ph, nullptr,
NUM_ELM*sizeof(TestType), true);
SECTION("H2H, H2PinMem and PinMem2H") {
HIP_CHECK(hipMemcpy(B_h, A_h, NUM_ELM*sizeof(TestType),
hipMemcpyDefault));
HIP_CHECK(hipMemcpy(A_Ph, B_h, NUM_ELM*sizeof(TestType),
hipMemcpyDefault));
HIP_CHECK(hipMemcpy(B_Ph, A_Ph, NUM_ELM*sizeof(TestType),
hipMemcpyDefault));
HipTest::checkTest(A_h, B_Ph, NUM_ELM);
}
SECTION("H2D-D2D-D2H-SameGPU") {
HIP_CHECK(hipMemcpy(A_d, A_h, NUM_ELM*sizeof(TestType), hipMemcpyDefault));
HIP_CHECK(hipMemcpy(B_d, A_d, NUM_ELM*sizeof(TestType), hipMemcpyDefault));
HIP_CHECK(hipMemcpy(B_h, B_d, NUM_ELM*sizeof(TestType), hipMemcpyDefault));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
SECTION("pH2D-D2D-D2pH-SameGPU") {
HIP_CHECK(hipMemcpy(A_d, A_Ph, NUM_ELM*sizeof(TestType),
hipMemcpyDefault));
HIP_CHECK(hipMemcpy(B_d, A_d, NUM_ELM*sizeof(TestType), hipMemcpyDefault));
HIP_CHECK(hipMemcpy(B_Ph, B_d, NUM_ELM*sizeof(TestType),
hipMemcpyDefault));
HipTest::checkTest(A_Ph, B_Ph, NUM_ELM);
}
SECTION("H2D-D2D-D2H-DeviceContextChange") {
int deviceCount = 0;
HIP_CHECK(hipGetDeviceCount(&deviceCount));
if (deviceCount < 2) {
SUCCEED("deviceCount less then 2");
} else {
int canAccessPeer = 0;
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipMemcpy(A_d, A_h, NUM_ELM*sizeof(TestType),
hipMemcpyDefault));
HIP_CHECK(hipMemcpy(B_d, A_d, NUM_ELM*sizeof(TestType),
hipMemcpyDefault));
HIP_CHECK(hipMemcpy(B_h, B_d, NUM_ELM*sizeof(TestType),
hipMemcpyDefault));
HipTest::checkTest(A_h, B_h, NUM_ELM);
} else {
SUCCEED("P2P capability is not present");
}
}
}
SECTION("H2D-D2D-D2H-PeerGPU") {
int deviceCount = 0;
HIP_CHECK(hipGetDeviceCount(&deviceCount));
if (deviceCount < 2) {
SUCCEED("deviceCount less then 2");
} else {
int canAccessPeer = 0;
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(1));
TestType *C_d{nullptr};
HipTest::initArrays<TestType>(nullptr, nullptr, &C_d,
nullptr, nullptr, nullptr,
NUM_ELM*sizeof(TestType));
HIP_CHECK(hipMemcpy(A_d, A_h, NUM_ELM*sizeof(TestType),
hipMemcpyDefault));
HIP_CHECK(hipMemcpy(C_d, A_d, NUM_ELM*sizeof(TestType),
hipMemcpyDefault));
HIP_CHECK(hipMemcpy(B_h, C_d, NUM_ELM*sizeof(TestType),
hipMemcpyDefault));
HipTest::checkTest(A_h, B_h, NUM_ELM);
HIP_CHECK(hipFree(C_d));
} else {
SUCCEED("P2P capability is not present");
}
}
}
HipTest::freeArrays<TestType>(A_d, B_d, nullptr, A_h, B_h, nullptr, false);
HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr, A_Ph,
B_Ph, nullptr, true);
}
/*
This testcase verfies the boundary checks of hipMemcpy API for different sizes
*/
TEST_CASE("Unit_hipMemcpy_BoundaryCheck") {
size_t maxElem = 32 * 1024 * 1024;
DeviceMemory<float> memD(maxElem);
HostMemory<float> memU(maxElem, 0 /*usePinnedHost*/);
HostMemory<float> memP(maxElem, 0 /*usePinnedHost*/);
memcpytest2<float>(&memD, &memU, 32 * 1024 * 1024, 0, 0, 0);
auto sizes = GENERATE(15 * 1024 * 1024, 16 * 1024 * 1024,
16 * 1024 * 1024 + 16 * 1024,
16 * 1024 * 1024 + 512 * 1024,
17 * 1024 * 1024 + 1024,
32 * 1024 * 1024);
memcpytest2<float>(&memD, &memP, sizes, 0, 0, 0);
}
/*
This testcase verifies the multi thread scenario
*/
TEST_CASE("Unit_hipMemcpy_MultiThreadWithSerialization") {
HIP_CHECK(hipDeviceReset());
// Simplest cases: serialize the threads, and also used pinned memory:
// This verifies that the sub-calls to memcpytest2 are correct.
multiThread_1<float>(true, true);
// Serialize, but use unpinned memory to stress the unpinned memory xfer path.
multiThread_1<float>(true, false);
}
/*
This testcase verifies the device offsets
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpy_DeviceOffsets", "", float, double) {
HIP_CHECK(hipDeviceReset());
size_t maxSize = 256 * 1024;
memcpytest2_offsets<TestType>(maxSize, true, false);
memcpytest2_offsets<TestType>(maxSize, false, true);
}
/*
This testcase verifies hipMemcpy API with pinnedMemory and hostRegister
along with kernel launches
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpy_PinnedRegMemWithKernelLaunch",
"", int, float, double) {
int numDevices = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices < 2) {
SUCCEED("No of devices are less than 2");
} else {
// 1 refers to pinned Memory
// 2 refers to register Memory
int MallocPinType = GENERATE(0, 1);
size_t Nbytes = NUM_ELM * sizeof(TestType);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU,
threadsPerBlock, NUM_ELM);
TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
TestType *X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr};
TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
if (MallocPinType) {
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, NUM_ELM, true);
} else {
A_h = reinterpret_cast<TestType*>(malloc(Nbytes));
HIP_CHECK(hipHostRegister(A_h, Nbytes, hipHostRegisterDefault));
B_h = reinterpret_cast<TestType*>(malloc(Nbytes));
HIP_CHECK(hipHostRegister(B_h, Nbytes, hipHostRegisterDefault));
C_h = reinterpret_cast<TestType*>(malloc(Nbytes));
HIP_CHECK(hipHostRegister(C_h, Nbytes, hipHostRegisterDefault));
HipTest::initArrays<TestType>(&A_d, &B_d, &C_d, nullptr, nullptr,
nullptr, NUM_ELM, false);
HipTest::setDefaultData<TestType>(NUM_ELM, A_h, B_h, C_h);
}
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, 0, static_cast<const TestType*>(A_d),
static_cast<const TestType*>(B_d), C_d, NUM_ELM);
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM);
unsigned int seed = time(0);
HIP_CHECK(hipSetDevice(rand_r(&seed) % (numDevices-1)+1));
int device;
hipGetDevice(&device);
std::cout <<"hipMemcpy is set to happen between device 0 and device "
<<device << std::endl;
HipTest::initArrays<TestType>(&X_d, &Y_d, &Z_d, nullptr,
nullptr, nullptr, NUM_ELM, false);
for (int j = 0; j < NUM_ELM; j++) {
A_h[j] = 0;
B_h[j] = 0;
C_h[j] = 0;
}
hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost);
hipMemcpy(X_d, A_h, Nbytes, hipMemcpyHostToDevice);
hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost);
hipMemcpy(Y_d, B_h, Nbytes, hipMemcpyHostToDevice);
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, 0, static_cast<const TestType*>(X_d),
static_cast<const TestType*>(Y_d), Z_d, NUM_ELM);
HIP_CHECK(hipMemcpy(C_h, Z_d, Nbytes, hipMemcpyDeviceToHost));
HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM);
if (MallocPinType) {
HipTest::freeArrays<TestType>(A_d, B_d, C_d, A_h, B_h, C_h, true);
} else {
HIP_CHECK(hipHostUnregister(A_h));
free(A_h);
HIP_CHECK(hipHostUnregister(B_h));
free(B_h);
HIP_CHECK(hipHostUnregister(C_h));
free(C_h);
HipTest::freeArrays<TestType>(A_d, B_d, C_d, nullptr,
nullptr, nullptr, false);
}
HipTest::freeArrays<TestType>(X_d, Y_d, Z_d, nullptr,
nullptr, nullptr, false);
}
}
catch/unit/memory/hipMemcpy2DFromArray.cc
+9 -4
View File
@@ -76,11 +76,14 @@ TEST_CASE("Unit_hipMemcpy2DFromArray_ExtentValidation") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
size_t width{sizeof(float)*NUM_W};
float *A_h{nullptr}, *hData{nullptr};
float *A_h{nullptr}, *hData{nullptr}, *valData{nullptr};
// Initialization of variables
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
&A_h, &hData, nullptr,
width*NUM_H, false);
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
nullptr, &valData, nullptr,
width*NUM_H, false);
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
@@ -101,9 +104,9 @@ TEST_CASE("Unit_hipMemcpy2DFromArray_ExtentValidation") {
A_h, width, width,
NUM_H, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d,
0, 0, NUM_W*sizeof(float),
0, 0, width,
0, hipMemcpyDeviceToHost));
REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) != true);
REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true);
}
// hipMemcpy2DFromArray API would return success for width and height as 0
// and does not perform any copy
@@ -120,13 +123,15 @@ TEST_CASE("Unit_hipMemcpy2DFromArray_ExtentValidation") {
HIP_CHECK(hipMemcpy2DFromArray(hData, width, A_d,
0, 0, 0,
NUM_H, hipMemcpyDeviceToHost));
REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) != true);
REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true);
}
// Cleaning the memory
HIP_CHECK(hipFreeArray(A_d));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, hData, nullptr, false);
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
nullptr, valData, nullptr, false);
}
/*
* This Scenario Verifies hipMemcpy2DFromArray API by copying the
catch/unit/memory/hipMemcpy2DFromArrayAsync.cc
+8 -3
View File
@@ -83,13 +83,16 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_ExtentValidation") {
HIP_CHECK(hipSetDevice(0));
hipArray *A_d{nullptr};
size_t width{sizeof(float)*NUM_W};
float *A_h{nullptr}, *hData{nullptr};
float *A_h{nullptr}, *hData{nullptr}, *valData{nullptr};
hipStream_t stream;
// Initialization of variables
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
&A_h, &hData, nullptr,
width*NUM_H, false);
HipTest::initArrays<float>(nullptr, nullptr, nullptr,
nullptr, &valData, nullptr,
width*NUM_H, false);
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
HIP_CHECK(hipMallocArray(&A_d, &desc, NUM_W, NUM_H, hipArrayDefault));
HIP_CHECK(hipStreamCreate(&stream));
@@ -116,7 +119,7 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_ExtentValidation") {
0, 0, NUM_W*sizeof(float),
0, hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) != true);
REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true);
}
// hipMemcpy2DFromArrayAsync API would return success for
// width and height as 0
@@ -135,7 +138,7 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_ExtentValidation") {
0, 0, 0,
NUM_H, hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
REQUIRE(HipTest::checkArray(hData, A_h, NUM_W, NUM_H) != true);
REQUIRE(HipTest::checkArray(hData, valData, NUM_W, NUM_H) == true);
}
// Cleaning the memory
@@ -143,6 +146,8 @@ TEST_CASE("Unit_hipMemcpy2DFromArrayAsync_ExtentValidation") {
HIP_CHECK(hipStreamDestroy(stream));
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
A_h, hData, nullptr, false);
HipTest::freeArrays<float>(nullptr, nullptr, nullptr,
nullptr, valData, nullptr, false);
}
/*
* This Scenario Verifies hipMemcpy2DFromArrayAsync API by copying the
catch/unit/memory/hipMemcpy3DAsync.cc
+7 -5
View File
@@ -479,11 +479,12 @@ void Memcpy3DAsync<T>::NegativeTests() {
template <typename T>
void Memcpy3DAsync<T>::D2D_SameDeviceMem_StreamDiffDevice() {
HIP_CHECK(hipSetDevice(0));
// Allocating the Memory
AllocateMemory();
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipStreamCreate(&stream));
SetDefaultData();
memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
SetDefaultData();
// Host to Device
myparms.srcPtr = make_hipPitchedPtr(hData, width * sizeof(T), width, height);
@@ -507,12 +508,13 @@ void Memcpy3DAsync<T>::D2D_SameDeviceMem_StreamDiffDevice() {
myparms.kind = hipMemcpyDeviceToDevice;
#endif
REQUIRE(hipMemcpy3DAsync(&myparms, stream) == hipSuccess);
memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
HIP_CHECK(hipStreamSynchronize(stream));
T *hOutputData = reinterpret_cast<T*>(malloc(size));
memset(hOutputData, 0, size);
SetDefaultData();
// Device to host
memset(&myparms, 0x0, sizeof(hipMemcpy3DParms));
SetDefaultData();
myparms.dstPtr = make_hipPitchedPtr(hOutputData,
width * sizeof(T), width, height);
myparms.srcArray = arr1;
@@ -727,8 +729,8 @@ TEST_CASE("Unit_hipMemcpy3DAsync_multiDevice-DiffStream") {
int numDevices = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
Memcpy3DAsync<int> memcpy3dAsync(width, height, depth,
hipChannelFormatKindSigned);
Memcpy3DAsync<float> memcpy3dAsync(width, height, depth,
hipChannelFormatKindFloat);
memcpy3dAsync.D2D_SameDeviceMem_StreamDiffDevice();
} else {
SUCCEED("skipping the testcases as numDevices < 2");
catch/unit/memory/hipMemcpyAllApiNegative.cc
+326
View File
@@ -0,0 +1,326 @@
/*
Copyright (c) 2021 - 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.
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
#include <utility>
#include <vector>
/*
This testfile verifies the following scenarios of all hipMemcpy API
1. Negative Scenarios
2. Half Memory copy scenarios
3. Null check scenario
*/
static constexpr auto NUM_ELM{1024*1024};
/*This testcase verifies the negative scenarios of hipMemcpy APIs
*/
TEST_CASE("Unit_hipMemcpy_Negative") {
// Initialization of variables
float *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
float *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
HIP_CHECK(hipSetDevice(0));
HipTest::initArrays<float>(&A_d, &B_d, &C_d,
&A_h, &B_h, &C_h,
NUM_ELM*sizeof(float));
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
SECTION("Pass nullptr to destination pointer for all Memcpy APIs") {
REQUIRE(hipMemcpy(nullptr, A_d, NUM_ELM * sizeof(float),
hipMemcpyDefault) != hipSuccess);
REQUIRE(hipMemcpyAsync(nullptr, A_h, NUM_ELM * sizeof(float),
hipMemcpyDefault, stream) != hipSuccess);
REQUIRE(hipMemcpyHtoD(hipDeviceptr_t(nullptr), A_h,
NUM_ELM * sizeof(float)) != hipSuccess);
REQUIRE(hipMemcpyHtoDAsync(hipDeviceptr_t(nullptr), A_h,
NUM_ELM * sizeof(float),
stream) != hipSuccess);
REQUIRE(hipMemcpyDtoH(nullptr, hipDeviceptr_t(A_d),
NUM_ELM * sizeof(float)) != hipSuccess);
REQUIRE(hipMemcpyDtoHAsync(nullptr, hipDeviceptr_t(A_d),
NUM_ELM * sizeof(float),
stream) != hipSuccess);
REQUIRE(hipMemcpyDtoD(hipDeviceptr_t(nullptr),
hipDeviceptr_t(A_d), NUM_ELM * sizeof(float))
!= hipSuccess);
REQUIRE(hipMemcpyDtoDAsync(hipDeviceptr_t(nullptr),
hipDeviceptr_t(A_d),
NUM_ELM * sizeof(float), stream)
!= hipSuccess);
}
SECTION("Passing nullptr to source pointer") {
REQUIRE(hipMemcpy(A_h, nullptr, NUM_ELM * sizeof(float),
hipMemcpyDefault) != hipSuccess);
REQUIRE(hipMemcpyAsync(A_d, nullptr,
NUM_ELM * sizeof(float),
hipMemcpyDefault, stream) != hipSuccess);
REQUIRE(hipMemcpyHtoD(hipDeviceptr_t(A_d), nullptr,
NUM_ELM * sizeof(float)) != hipSuccess);
REQUIRE(hipMemcpyHtoDAsync(hipDeviceptr_t(A_d), nullptr,
NUM_ELM * sizeof(float),
stream) != hipSuccess);
REQUIRE(hipMemcpyDtoH(A_h, hipDeviceptr_t(nullptr),
NUM_ELM * sizeof(float)) != hipSuccess);
REQUIRE(hipMemcpyDtoHAsync(A_h, hipDeviceptr_t(nullptr),
NUM_ELM * sizeof(float),
stream) != hipSuccess);
REQUIRE(hipMemcpyDtoD(hipDeviceptr_t(A_d),
hipDeviceptr_t(nullptr), NUM_ELM * sizeof(float))
!= hipSuccess);
REQUIRE(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d),
hipDeviceptr_t(nullptr),
NUM_ELM * sizeof(float), stream)
!= hipSuccess);
}
SECTION("Passing nullptr to both source and dest pointer") {
REQUIRE(hipMemcpy(nullptr, nullptr, NUM_ELM * sizeof(float),
hipMemcpyDefault) != hipSuccess);
REQUIRE(hipMemcpyAsync(nullptr, nullptr, NUM_ELM * sizeof(float),
hipMemcpyDefault, stream) != hipSuccess);
REQUIRE(hipMemcpyHtoD(hipDeviceptr_t(nullptr), nullptr,
NUM_ELM * sizeof(float)) != hipSuccess);
REQUIRE(hipMemcpyHtoDAsync(hipDeviceptr_t(nullptr), nullptr,
NUM_ELM * sizeof(float),
stream) != hipSuccess);
REQUIRE(hipMemcpyDtoH(nullptr, hipDeviceptr_t(nullptr),
NUM_ELM * sizeof(float)) != hipSuccess);
REQUIRE(hipMemcpyDtoHAsync(nullptr, hipDeviceptr_t(nullptr),
NUM_ELM * sizeof(float),
stream) != hipSuccess);
REQUIRE(hipMemcpyDtoD(hipDeviceptr_t(nullptr),
hipDeviceptr_t(nullptr), NUM_ELM * sizeof(float))
!= hipSuccess);
REQUIRE(hipMemcpyDtoDAsync(hipDeviceptr_t(nullptr),
hipDeviceptr_t(nullptr),
NUM_ELM * sizeof(float), stream)
!= hipSuccess);
}
SECTION("Passing same pointers") {
HIP_CHECK(hipMemcpy(A_d, A_d, (NUM_ELM/2) * sizeof(float),
hipMemcpyDefault));
HIP_CHECK(hipMemcpy(A_h, A_h, (NUM_ELM/2) * sizeof(float),
hipMemcpyDefault));
HIP_CHECK(hipMemcpyAsync(A_d, A_d, (NUM_ELM/2) * sizeof(float),
hipMemcpyDefault, stream));
HIP_CHECK(hipMemcpyAsync(A_h, A_h, (NUM_ELM/2) * sizeof(float),
hipMemcpyDefault, stream));
HIP_CHECK(hipMemcpyDtoD(hipDeviceptr_t(A_d),
hipDeviceptr_t(A_d),
NUM_ELM * sizeof(float)));
HIP_CHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d),
hipDeviceptr_t(A_d),
NUM_ELM * sizeof(float), stream));
}
HipTest::freeArrays<float>(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipStreamDestroy(stream));
}
/*
This testcase verifies the Nullcheck for all the 8 Memcpy APIs
*/
TEST_CASE("Unit_hipMemcpy_NullCheck") {
// Initialization of variables
float *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
float *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
HIP_CHECK(hipSetDevice(0));
HipTest::initArrays<float>(&A_d, &B_d, &C_d,
&A_h, &B_h, &C_h,
NUM_ELM*sizeof(float));
hipStream_t stream;
hipStreamCreate(&stream);
HIP_CHECK(hipMemcpy(A_d, C_h,
NUM_ELM*sizeof(float),
hipMemcpyHostToDevice));
SECTION("hipMemcpyHtoD API null size check") {
REQUIRE(hipMemcpyHtoD(hipDeviceptr_t(A_d), A_h,
0) == hipSuccess);
HIP_CHECK(hipMemcpy(B_h, A_d,
NUM_ELM*sizeof(float),
hipMemcpyDeviceToHost));
HipTest::checkTest(C_h, B_h, NUM_ELM);
}
SECTION("hipMemcpyHtoDAsync API null size check") {
HIP_CHECK(hipMemcpyHtoDAsync(hipDeviceptr_t(A_d), A_h,
0, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipMemcpy(B_h, A_d,
NUM_ELM*sizeof(float),
hipMemcpyDeviceToHost));
HipTest::checkTest(C_h, B_h, NUM_ELM);
}
SECTION("hipMemcpy API null size check") {
HIP_CHECK(hipMemcpy(A_d, B_h, 0, hipMemcpyDefault));
HIP_CHECK(hipMemcpy(B_h, A_d,
NUM_ELM*sizeof(float),
hipMemcpyDeviceToHost));
HipTest::checkTest(C_h, B_h, NUM_ELM);
}
SECTION("hipMemcpyAsync API null size check") {
HIP_CHECK(hipMemcpyAsync(A_d, B_h, 0, hipMemcpyDefault, stream));
HIP_CHECK(hipMemcpy(B_h, A_d,
NUM_ELM*sizeof(float),
hipMemcpyDeviceToHost));
HipTest::checkTest(C_h, B_h, NUM_ELM);
}
SECTION("hipMemcpyDtoH API null size check") {
HIP_CHECK(hipMemcpyDtoH(C_h, hipDeviceptr_t(A_d), 0));
HIP_CHECK(hipMemcpy(B_h, A_d,
NUM_ELM*sizeof(float),
hipMemcpyDeviceToHost));
HipTest::checkTest(C_h, B_h, NUM_ELM);
}
SECTION("hipMemcpyDtoHAsync API null size check") {
HIP_CHECK(hipMemcpyDtoHAsync(C_h, hipDeviceptr_t(A_d), 0, stream));
HIP_CHECK(hipMemcpy(B_h, A_d,
NUM_ELM*sizeof(float),
hipMemcpyDeviceToHost));
HipTest::checkTest(C_h, B_h, NUM_ELM);
}
SECTION("hipMemcpyDtoD API null size check") {
HIP_CHECK(hipMemcpy(C_d, A_h,
NUM_ELM*sizeof(float),
hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpyDtoD(hipDeviceptr_t(C_d), hipDeviceptr_t(A_d), 0));
HIP_CHECK(hipMemcpy(B_h, C_d,
NUM_ELM*sizeof(float),
hipMemcpyDeviceToHost));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
SECTION("hipMemcpyDtoDAsync API null size check") {
HIP_CHECK(hipMemcpy(C_d, A_h,
NUM_ELM*sizeof(float),
hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(C_d), hipDeviceptr_t(A_d),
0, stream));
HIP_CHECK(hipMemcpy(B_h, C_d,
NUM_ELM*sizeof(float),
hipMemcpyDeviceToHost));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
HipTest::freeArrays<float>(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipStreamDestroy(stream));
}
/*
This testcase verifies all the hipMemcpy APIs by
copying half the memory.
*/
TEST_CASE("Unit_hipMemcpy_HalfMemCopy") {
// Initialization of variables
float *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
float *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
HIP_CHECK(hipSetDevice(0));
HipTest::initArrays<float>(&A_d, &B_d, &C_d,
&A_h, &B_h, &C_h,
NUM_ELM*sizeof(float));
hipStream_t stream;
hipStreamCreate(&stream);
SECTION("hipMemcpyHtoD half memory copy") {
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d), A_h,
(NUM_ELM * sizeof(float))/2));
HIP_CHECK(hipMemcpy(B_h, A_d,
(NUM_ELM*sizeof(float))/2,
hipMemcpyDeviceToHost));
HipTest::checkTest(A_h, B_h, NUM_ELM/2);
}
SECTION("hipMemcpyHtoDAsync half memory copy") {
HIP_CHECK(hipMemcpyHtoDAsync(hipDeviceptr_t(A_d), A_h,
(NUM_ELM * sizeof(float))/2, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipMemcpy(B_h, A_d,
(NUM_ELM*sizeof(float))/2,
hipMemcpyDeviceToHost));
HipTest::checkTest(A_h, B_h, NUM_ELM/2);
}
SECTION("hipMemcpyDtoH half memory copy") {
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d), A_h,
(NUM_ELM * sizeof(float))));
HIP_CHECK(hipMemcpyDtoH(B_h, hipDeviceptr_t(A_d),
(NUM_ELM * sizeof(float))/2));
HipTest::checkTest(A_h, B_h, NUM_ELM/2);
}
SECTION("hipMemcpyDtoHAsync half memory copy") {
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d), A_h,
(NUM_ELM * sizeof(float))));
HIP_CHECK(hipMemcpyDtoHAsync(B_h, hipDeviceptr_t(A_d),
(NUM_ELM * sizeof(float))/2,
stream));
HIP_CHECK(hipStreamSynchronize(stream));
HipTest::checkTest(A_h, B_h, NUM_ELM/2);
}
SECTION("hipMemcpyDtoD half memory copy") {
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d), A_h,
(NUM_ELM * sizeof(float))/2));
HIP_CHECK(hipMemcpyDtoD(hipDeviceptr_t(B_d), hipDeviceptr_t(A_d),
(NUM_ELM*sizeof(float))/2));
HIP_CHECK(hipMemcpy(B_h, B_d,
(NUM_ELM*sizeof(float))/2,
hipMemcpyDeviceToHost));
HipTest::checkTest(A_h, B_h, NUM_ELM/2);
}
SECTION("hipMemcpyDtoDAsync half memory copy") {
HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d), A_h,
(NUM_ELM * sizeof(float))/2));
HIP_CHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(B_d), hipDeviceptr_t(A_d),
(NUM_ELM*sizeof(float))/2,
stream));
HIP_CHECK(hipMemcpy(B_h, B_d,
(NUM_ELM*sizeof(float))/2,
hipMemcpyDeviceToHost));
HipTest::checkTest(A_h, B_h, NUM_ELM/2);
}
SECTION("hipMemcpy half memory copy") {
HIP_CHECK(hipMemcpy(A_d, A_h
, (NUM_ELM*sizeof(float)),
hipMemcpyDefault));
HIP_CHECK(hipMemcpy(B_h, A_d,
(NUM_ELM/2)*sizeof(float),
hipMemcpyDeviceToHost));
HipTest::checkTest(A_h, B_h, NUM_ELM/2);
}
SECTION("hipMemcpyAsync half memory copy") {
HIP_CHECK(hipMemcpy(A_d, A_h,
(NUM_ELM*sizeof(float)),
hipMemcpyDefault));
HIP_CHECK(hipMemcpyAsync(B_h, A_d,
(NUM_ELM/2)*sizeof(float),
hipMemcpyDeviceToHost, stream));
HipTest::checkTest(A_h, B_h, NUM_ELM/2);
}
HipTest::freeArrays<float>(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipStreamDestroy(stream));
}
catch/unit/memory/hipMemcpyAsync.cc
+403
View File
@@ -0,0 +1,403 @@
/*
Copyright (c) 2021 - 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.
*/
/*
This testcase verifies the following scenarios
1. hipMemcpyAsync with kernel launch
2. H2D-D2D-D2H-H2PinnMem and device context change scenarios
3. This test launches multiple threads which uses same stream to deploy kernel
and also launch hipMemcpyAsync() api. This test case is simulate the scenario
reported in SWDEV-181598.
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
#include <atomic>
#define NUM_THREADS 16
static constexpr auto NUM_ELM{1024 * 1024};
static constexpr size_t N_ELMTS{32 * 1024};
std::atomic<size_t> Thread_count { 0 };
static unsigned blocksPerCU{6}; // to hide latency
static unsigned threadsPerBlock{256};
template<typename T>
void Thread_func(T *A_d, T *B_d, T* C_d, T* C_h, size_t Nbytes,
hipStream_t mystream) {
unsigned blocks = HipTest::setNumBlocks(blocksPerCU,
threadsPerBlock, N_ELMTS);
hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
dim3(threadsPerBlock), 0,
mystream, A_d, C_d, N_ELMTS);
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream));
// The following two MemcpyAsync calls are for sole
// purpose of loading stream with multiple async calls
HIP_CHECK(hipMemcpyAsync(B_d, A_d, Nbytes,
hipMemcpyDeviceToDevice, mystream));
HIP_CHECK(hipMemcpyAsync(B_d, A_d, Nbytes,
hipMemcpyDeviceToDevice, mystream));
Thread_count++;
}
template<typename T>
void Thread_func_MultiStream() {
int Data_mismatch = 0;
T *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
T *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
size_t Nbytes = N_ELMTS * sizeof(T);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU,
threadsPerBlock, N_ELMTS);
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N_ELMTS, false);
hipStream_t mystream;
HIP_CHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking));
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream));
hipLaunchKernelGGL((HipTest::vector_square), dim3(blocks),
dim3(threadsPerBlock), 0,
mystream, A_d, C_d, N_ELMTS);
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream));
// The following hipMemcpyAsync() is called only to
// load stream with multiple Async calls
HIP_CHECK(hipMemcpyAsync(B_d, A_d, Nbytes,
hipMemcpyDeviceToDevice, mystream));
Thread_count++;
HIP_CHECK(hipStreamSynchronize(mystream));
HIP_CHECK(hipStreamDestroy(mystream));
// Verifying result of the kernel computation
for (size_t i = 0; i < N_ELMTS; i++) {
if (C_h[i] != A_h[i] * A_h[i]) {
Data_mismatch++;
}
}
// Releasing resources
HipTest::freeArrays<T>(A_d, B_d, C_d, A_h, B_h, C_h, false);
REQUIRE(Data_mismatch == 0);
}
/*
This testcase verifies hipMemcpyAsync API
Initializes device variables
Launches kernel and performs the sum of device variables
copies the result to host variable and validates the result.
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_KernelLaunch", "", int, float,
double) {
size_t Nbytes = NUM_ELM * sizeof(TestType);
TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
HIP_CHECK(hipSetDevice(0));
hipStream_t stream;
hipStreamCreate(&stream);
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, NUM_ELM, false);
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream));
HIP_CHECK(hipStreamSynchronize(stream));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0,
static_cast<const TestType*>(A_d),
static_cast<const TestType*>(B_d), C_d, NUM_ELM);
HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM);
HipTest::freeArrays<TestType>(A_d, B_d, C_d, A_h, B_h, C_h, false);
}
/*
This testcase verifies the following scenarios
1. H2H,H2PinMem and PinnedMem2Host
2. H2D-D2D-D2H in same GPU
3. Pinned Host Memory to device variables in same GPU
4. Device context change
5. H2D-D2D-D2H peer GPU
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_H2H-H2D-D2H-H2PinMem", "", char, int,
float, double) {
TestType *A_d{nullptr}, *B_d{nullptr};
TestType *A_h{nullptr}, *B_h{nullptr};
TestType *A_Ph{nullptr}, *B_Ph{nullptr};
HIP_CHECK(hipSetDevice(0));
hipStream_t stream;
hipStreamCreate(&stream);
HipTest::initArrays<TestType>(&A_d, &B_d, nullptr,
&A_h, &B_h, nullptr,
NUM_ELM*sizeof(TestType));
HipTest::initArrays<TestType>(nullptr, nullptr, nullptr,
&A_Ph, &B_Ph, nullptr,
NUM_ELM*sizeof(TestType), true);
SECTION("H2H, H2PinMem and PinMem2H") {
HIP_CHECK(hipMemcpyAsync(B_h, A_h, NUM_ELM*sizeof(TestType),
hipMemcpyHostToHost, stream));
HIP_CHECK(hipMemcpyAsync(A_Ph, B_h, NUM_ELM*sizeof(TestType),
hipMemcpyHostToHost, stream));
HIP_CHECK(hipMemcpyAsync(B_Ph, A_Ph, NUM_ELM*sizeof(TestType),
hipMemcpyHostToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HipTest::checkTest(A_h, B_Ph, NUM_ELM);
}
SECTION("H2D-D2D-D2H-SameGPU") {
HIP_CHECK(hipMemcpyAsync(A_d, A_h, NUM_ELM*sizeof(TestType),
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemcpyAsync(B_d, A_d, NUM_ELM*sizeof(TestType),
hipMemcpyDeviceToDevice, stream));
HIP_CHECK(hipMemcpyAsync(B_h, B_d, NUM_ELM*sizeof(TestType),
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HipTest::checkTest(A_h, B_h, NUM_ELM);
}
SECTION("pH2D-D2D-D2pH-SameGPU") {
HIP_CHECK(hipMemcpyAsync(A_d, A_Ph, NUM_ELM*sizeof(TestType),
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemcpyAsync(B_d, A_d, NUM_ELM*sizeof(TestType),
hipMemcpyDeviceToDevice, stream));
HIP_CHECK(hipMemcpyAsync(B_Ph, B_d, NUM_ELM*sizeof(TestType),
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HipTest::checkTest(A_Ph, B_Ph, NUM_ELM);
}
SECTION("H2D-D2D-D2H-DeviceContextChange") {
int deviceCount = 0;
HIP_CHECK(hipGetDeviceCount(&deviceCount));
if (deviceCount < 2) {
SUCCEED("deviceCount less then 2");
} else {
int canAccessPeer = 0;
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipMemcpyAsync(A_d, A_h, NUM_ELM*sizeof(TestType),
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemcpyAsync(B_d, A_d, NUM_ELM*sizeof(TestType),
hipMemcpyDeviceToDevice, stream));
HIP_CHECK(hipMemcpyAsync(B_h, B_d, NUM_ELM*sizeof(TestType),
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HipTest::checkTest(A_h, B_h, NUM_ELM);
} else {
SUCCEED("P2P capability is not present");
}
}
}
SECTION("H2D-D2D-D2H-PeerGPU") {
int deviceCount = 0;
HIP_CHECK(hipGetDeviceCount(&deviceCount));
if (deviceCount < 2) {
SUCCEED("deviceCount less then 2");
} else {
int canAccessPeer = 0;
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(1));
TestType *C_d{nullptr};
HipTest::initArrays<TestType>(nullptr, nullptr, &C_d,
nullptr, nullptr, nullptr,
NUM_ELM*sizeof(TestType));
HIP_CHECK(hipMemcpyAsync(A_d, A_h, NUM_ELM*sizeof(TestType),
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemcpyAsync(C_d, A_d, NUM_ELM*sizeof(TestType),
hipMemcpyDeviceToDevice, stream));
HIP_CHECK(hipMemcpyAsync(B_h, C_d, NUM_ELM*sizeof(TestType),
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HipTest::checkTest(A_h, B_h, NUM_ELM);
HIP_CHECK(hipFree(C_d));
} else {
SUCCEED("P2P capability is not present");
}
}
}
HipTest::freeArrays<TestType>(A_d, B_d, nullptr, A_h, B_h, nullptr, false);
HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr, A_Ph,
B_Ph, nullptr, true);
}
// This test launches multiple threads which uses same stream to deploy kernel
// and also launch hipMemcpyAsync() api. This test case is simulate the scenario
// reported in SWDEV-181598
TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_hipMultiMemcpyMultiThread", "",
int, float, double) {
size_t Nbytes = N_ELMTS * sizeof(TestType);
int Data_mismatch = 0;
hipStream_t mystream;
TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N_ELMTS, false);
HIP_CHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking));
HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream));
std::thread T[NUM_THREADS];
for (int i = 0; i < NUM_THREADS; i++) {
T[i] = std::thread(Thread_func<TestType>, A_d, B_d, C_d,
C_h, Nbytes, mystream);
}
// Wait until all the threads finish their execution
for (int i = 0; i < NUM_THREADS; i++) {
T[i].join();
}
HIP_CHECK(hipStreamSynchronize(mystream));
HIP_CHECK(hipStreamDestroy(mystream));
// Verifying the result of the kernel computation
for (size_t i = 0; i < N_ELMTS; i++) {
if (C_h[i] != A_h[i] * A_h[i]) {
Data_mismatch++;
}
}
REQUIRE(Thread_count.load() == NUM_THREADS);
REQUIRE(Data_mismatch == 0);
HipTest::freeArrays<TestType>(A_d, B_d, C_d, A_h, B_h, C_h, false);
Thread_count.exchange(0);
}
TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_hipMultiMemcpyMultiThreadMultiStream",
"", int, float, double) {
std::thread T[NUM_THREADS];
for (int i = 0; i < NUM_THREADS; i++) {
T[i] = std::thread(Thread_func_MultiStream<TestType>);
}
// Wait until all the threads finish their execution
for (int i = 0; i < NUM_THREADS; i++) {
T[i].join();
}
REQUIRE(Thread_count.load() == NUM_THREADS);
Thread_count.exchange(0);
}
/*
This testcase verifies hipMemcpy API with pinnedMemory and hostRegister
along with kernel launches
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_PinnedRegMemWithKernelLaunch",
"", int, float, double) {
int numDevices = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices < 2) {
SUCCEED("No of devices are less than 2");
} else {
// 1 refers to pinned Memory
// 2 refers to register Memory
int MallocPinType = GENERATE(0, 1);
size_t Nbytes = NUM_ELM * sizeof(TestType);
unsigned blocks = HipTest::setNumBlocks(blocksPerCU,
threadsPerBlock, NUM_ELM);
TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr};
TestType *X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr};
TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
if (MallocPinType) {
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, NUM_ELM, true);
} else {
A_h = reinterpret_cast<TestType*>(malloc(Nbytes));
HIP_CHECK(hipHostRegister(A_h, Nbytes, hipHostRegisterDefault));
B_h = reinterpret_cast<TestType*>(malloc(Nbytes));
HIP_CHECK(hipHostRegister(B_h, Nbytes, hipHostRegisterDefault));
C_h = reinterpret_cast<TestType*>(malloc(Nbytes));
HIP_CHECK(hipHostRegister(C_h, Nbytes, hipHostRegisterDefault));
HipTest::initArrays<TestType>(&A_d, &B_d, &C_d, nullptr, nullptr,
nullptr, NUM_ELM, false);
HipTest::setDefaultData<TestType>(NUM_ELM, A_h, B_h, C_h);
}
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, 0, static_cast<const TestType*>(A_d),
static_cast<const TestType*>(B_d), C_d, NUM_ELM);
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM);
unsigned int seed = time(0);
HIP_CHECK(hipSetDevice(rand_r(&seed) % (numDevices-1)+1));
int device;
HIP_CHECK(hipGetDevice(&device));
INFO("hipMemcpy is set to happen between device 0 and device "
<< device);
HipTest::initArrays<TestType>(&X_d, &Y_d, &Z_d, nullptr,
nullptr, nullptr, NUM_ELM, false);
hipStream_t gpu1Stream;
HIP_CHECK(hipStreamCreate(&gpu1Stream));
for (int j = 0; j < NUM_ELM; j++) {
A_h[j] = 0;
B_h[j] = 0;
C_h[j] = 0;
}
hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost);
hipMemcpyAsync(X_d, A_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream);
hipMemcpy(B_h, B_d, Nbytes, hipMemcpyDeviceToHost);
hipMemcpyAsync(Y_d, B_h, Nbytes, hipMemcpyHostToDevice, gpu1Stream);
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, 0, static_cast<const TestType*>(X_d),
static_cast<const TestType*>(Y_d), Z_d, NUM_ELM);
HIP_CHECK(hipMemcpyAsync(C_h, Z_d, Nbytes,
hipMemcpyDeviceToHost, gpu1Stream));
HIP_CHECK(hipStreamSynchronize(gpu1Stream));
HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM);
if (MallocPinType) {
HipTest::freeArrays<TestType>(A_d, B_d, C_d, A_h, B_h, C_h, true);
} else {
HIP_CHECK(hipHostUnregister(A_h));
free(A_h);
HIP_CHECK(hipHostUnregister(B_h));
free(B_h);
HIP_CHECK(hipHostUnregister(C_h));
free(C_h);
HipTest::freeArrays<TestType>(A_d, B_d, C_d, nullptr,
nullptr, nullptr, false);
}
HipTest::freeArrays<TestType>(X_d, Y_d, Z_d, nullptr,
nullptr, nullptr, false);
HIP_CHECK(hipStreamDestroy(gpu1Stream));
}
}
catch/unit/memory/hipMemcpyDtoD.cc
+95
View File
@@ -0,0 +1,95 @@
/*
Copyright (c) 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/*
This testcase verifies the hipMemcpyDtoD basic scenario
1. H2D-KernelLaunch-D2H scenario
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
static constexpr auto NUM_ELM{1024};
/*
This testcase verifies hipMemcpyDtoD API
1.Initializes device variables
2.Launches kernel and performs the sum of device variables
3.Copies the result to host variable and validates the result.
4.Sets the peer device
5.D2D copy from GPU-0 to GPU-1
6.Kernel Launch
7.DtoH copy and validating the result
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpyDtoD_Basic", "",
int, float, double) {
size_t Nbytes = NUM_ELM * sizeof(TestType);
int numDevices = 0;
TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr},
*X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr};
TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int canAccessPeer = 0;
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(0));
HipTest::initArrays<TestType>(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h,
NUM_ELM, false);
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipMalloc(&X_d, Nbytes));
HIP_CHECK(hipMalloc(&Y_d, Nbytes));
HIP_CHECK(hipMalloc(&Z_d, Nbytes));
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1),
dim3(1), 0, 0,
static_cast<const TestType *>(A_d),
static_cast<const TestType *>(B_d), C_d, NUM_ELM);
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipDeviceSynchronize());
HipTest::checkVectorADD<TestType>(A_h, B_h, C_h, NUM_ELM);
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipMemcpyDtoD((hipDeviceptr_t)X_d, (hipDeviceptr_t)A_d,
Nbytes));
HIP_CHECK(hipMemcpyDtoD((hipDeviceptr_t)Y_d, (hipDeviceptr_t)B_d,
Nbytes));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1),
dim3(1), 0, 0,
static_cast<const TestType*>(X_d),
static_cast<const TestType*>(Y_d), Z_d, NUM_ELM);
HIP_CHECK(hipMemcpyDtoH(C_h, (hipDeviceptr_t)Z_d, Nbytes));
HIP_CHECK(hipDeviceSynchronize());
HipTest::checkVectorADD<TestType>(A_h, B_h, C_h, NUM_ELM);
HipTest::freeArrays<TestType>(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipFree(X_d));
HIP_CHECK(hipFree(Y_d));
HIP_CHECK(hipFree(Z_d));
} else {
SUCCEED("Machine does not seem to have P2P Capabilities");
}
}
}
catch/unit/memory/hipMemcpyDtoDAsync.cc
+100
View File
@@ -0,0 +1,100 @@
/*
Copyright (c) 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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/*
This testcase verifies the Basic scenario
1. H2D-KernelLaunch-D2H then D2D-KernelLaunch-D2H in peer GPU
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
static constexpr auto NUM_ELM{1024};
/*
This testcase verifies hipMemcpyDtoDAsync API
1.Initializes device variables
2.Launches kernel and performs the sum of device variables
3.Copies the result to host variable and validates the result.
4.Sets the peer device
5.D2D copy from GPU-0 to GPU-1
6.Kernel Launch
7.DtoH copy and validating the result
*/
TEMPLATE_TEST_CASE("Unit_hipMemcpyDtoDAsync_Basic", "",
int, float, double) {
size_t Nbytes = NUM_ELM * sizeof(TestType);
int numDevices = 0;
TestType *A_d{nullptr}, *B_d{nullptr}, *C_d{nullptr},
*X_d{nullptr}, *Y_d{nullptr}, *Z_d{nullptr};
TestType *A_h{nullptr}, *B_h{nullptr}, *C_h{nullptr};
hipStream_t stream;
HIP_CHECK(hipGetDeviceCount(&numDevices));
if (numDevices > 1) {
int canAccessPeer = 0;
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1));
if (canAccessPeer) {
HIP_CHECK(hipSetDevice(0));
HipTest::initArrays<TestType>(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h,
NUM_ELM, false);
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipMalloc(&X_d, Nbytes));
HIP_CHECK(hipMalloc(&Y_d, Nbytes));
HIP_CHECK(hipMalloc(&Z_d, Nbytes));
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1),
dim3(1), 0, 0,
static_cast<const TestType *>(A_d),
static_cast<const TestType *>(B_d), C_d, NUM_ELM);
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HIP_CHECK(hipDeviceSynchronize());
HipTest::checkVectorADD<TestType>(A_h, B_h, C_h, NUM_ELM);
HIP_CHECK(hipSetDevice(1));
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemcpyDtoDAsync((hipDeviceptr_t)X_d, (hipDeviceptr_t)A_d,
Nbytes, stream));
HIP_CHECK(hipMemcpyDtoDAsync((hipDeviceptr_t)Y_d, (hipDeviceptr_t)B_d,
Nbytes, stream));
HIP_CHECK(hipStreamSynchronize(stream));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(1),
dim3(1), 0, 0,
static_cast<const TestType*>(X_d),
static_cast<const TestType*>(Y_d), Z_d, NUM_ELM);
HIP_CHECK(hipMemcpyDtoHAsync(C_h, (hipDeviceptr_t)Z_d, Nbytes, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipDeviceSynchronize());
HipTest::checkVectorADD<TestType>(A_h, B_h, C_h, NUM_ELM);
HipTest::freeArrays<TestType>(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipFree(X_d));
HIP_CHECK(hipFree(Y_d));
HIP_CHECK(hipFree(Z_d));
} else {
SUCCEED("Machine does not seem to have P2P Capabilities");
}
}
}
catch/unit/memory/hipMemcpyWithStream.cc
+628
View File
@@ -0,0 +1,628 @@
/*
Copyright (c) 2021-22-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.
*/
/*
* Different test for checking functionality of
* hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,hipMemcpyKind kind,
* hipStream_t stream);
*/
/*
This testfile verifies the following scenarios
1. hipMemcpyWithStream with one stream
2. hipMemcpyWithStream with two streams
3. Multi GPU and single stream
4. hipMemcpyWithStream API with testkind DtoH
5. hipMemcpyWithStream API with testkind DtoD
6. hipMemcpyWithStream API with testkind HtoH
7. hipMemcpyWithStream API with testkind TestkindDefault
8. hipMemcpyWithStream API with testkind TestkindDefaultForDtoD
9. hipMemcpyWithStream API DtoD on same device
10.Multi threaded scenario
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
#include<vector>
#include<thread>
#include<chrono>
#define LEN 64
#define SIZE LEN << 2
#define THREADS 2
#define MAX_THREADS 16
static constexpr size_t N{4 * 1024 * 1024};
static const auto MaxGPUDevices{256};
static constexpr unsigned blocksPerCU{6}; // to hide latency
static constexpr unsigned threadsPerBlock{256};
enum class ops
{ TestwithOnestream,
TestwithTwoStream,
TestOnMultiGPUwithOneStream,
TestkindDtoH,
TestkindDtoD,
TestkindHtoH,
TestkindDefault,
TestkindDefaultForDtoD,
TestDtoDonSameDevice,
END_OF_LIST
};
struct joinable_thread : std::thread {
template <class... Xs>
explicit joinable_thread(Xs&&... xs) : std::thread(std::forward<Xs>(xs)...)
{} // NOLINT
joinable_thread& operator=(joinable_thread&& other) = default;
joinable_thread(joinable_thread&& other) = default;
~joinable_thread() {
if (this->joinable())
this->join();
}
};
void TestwithOnestream(void) {
size_t Nbytes = N * sizeof(int);
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemcpyWithStream(A_d, A_h, Nbytes,
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemcpyWithStream(B_d, B_h, Nbytes,
hipMemcpyHostToDevice, stream));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
HipTest::checkVectorADD(A_h, B_h, C_h, N);
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipStreamDestroy(stream));
}
void TestwithTwoStream(void) {
size_t Nbytes = N * sizeof(int);
const int NUM_STREAMS = 2;
int *A_d[NUM_STREAMS], *B_d[NUM_STREAMS], *C_d[NUM_STREAMS];
int *A_h[NUM_STREAMS], *B_h[NUM_STREAMS], *C_h[NUM_STREAMS];
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
for (int i=0; i < NUM_STREAMS; ++i) {
HipTest::initArrays(&A_d[i], &B_d[i], &C_d[i],
&A_h[i], &B_h[i], &C_h[i], N, false);
}
hipStream_t stream[NUM_STREAMS];
for (int i=0; i < NUM_STREAMS; ++i) {
HIP_CHECK(hipStreamCreate(&stream[i]));
}
for (int i=0; i < NUM_STREAMS; ++i) {
HIP_CHECK(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes,
hipMemcpyHostToDevice, stream[i]));
HIP_CHECK(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes,
hipMemcpyHostToDevice, stream[i]));
}
for (int i=0; i < NUM_STREAMS; ++i) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
}
for (int i=0; i < NUM_STREAMS; ++i) {
HIP_CHECK(hipStreamSynchronize(stream[i]));
HIP_CHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
HipTest::checkVectorADD(A_h[i], B_h[i], C_h[i], N);
}
for (int i=0; i < NUM_STREAMS; ++i) {
HipTest::freeArrays(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i], false);
HIP_CHECK(hipStreamDestroy(stream[i]));
}
}
void TestDtoDonSameDevice(void) {
size_t Nbytes = N * sizeof(int);
const int NUM_STREAMS = 2;
int *A_d[NUM_STREAMS], *B_d[NUM_STREAMS], *C_d[NUM_STREAMS];
int *A_h[NUM_STREAMS], *B_h[NUM_STREAMS], *C_h[NUM_STREAMS];
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0],
&A_h[0], &B_h[0], &C_h[0], N, false);
hipStream_t stream[NUM_STREAMS];
for (int i=0; i < NUM_STREAMS; ++i) {
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipStreamCreate(&stream[i]));
}
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipMalloc(&A_d[1], Nbytes));
HIP_CHECK(hipMalloc(&B_d[1], Nbytes));
HIP_CHECK(hipMalloc(&C_d[1], Nbytes));
C_h[1] = reinterpret_cast<int*>(malloc(Nbytes));
HIP_ASSERT(C_h[1] != NULL);
HIP_CHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes,
hipMemcpyHostToDevice, stream[0]));
HIP_CHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes,
hipMemcpyHostToDevice, stream[0]));
HIP_CHECK(hipMemcpyWithStream(A_d[1], A_d[0], Nbytes,
hipMemcpyDeviceToDevice, stream[1]));
HIP_CHECK(hipMemcpyWithStream(B_d[1], B_d[0], Nbytes,
hipMemcpyDeviceToDevice, stream[1]));
for (int i=0; i < NUM_STREAMS; ++i) {
HIP_CHECK(hipSetDevice(0));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
}
for (int i=0; i < NUM_STREAMS; ++i) {
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipStreamSynchronize(stream[i]));
HIP_CHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N);
}
HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false);
if (A_d[1]) {
HIP_CHECK(hipFree(A_d[1]));
}
if (B_d[1]) {
HIP_CHECK(hipFree(B_d[1]));
}
if (C_d[1]) {
HIP_CHECK(hipFree(C_d[1]));
}
if (C_h[1]) {
free(C_h[1]);
}
for (int i=0; i < NUM_STREAMS; ++i) {
HIP_CHECK(hipStreamDestroy(stream[i]));
}
}
void TestOnMultiGPUwithOneStream(void) {
size_t Nbytes = N * sizeof(int);
int NumDevices = 0;
HIP_CHECK(hipGetDeviceCount(&NumDevices));
// If you have single GPU machine the return
if (NumDevices <= 1) {
SUCCEED("NumDevices <2");
} else {
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices];
int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices];
hipStream_t stream[MaxGPUDevices];
for (int i=0; i < NumDevices; ++i) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipStreamCreate(&stream[i]));
}
for (int i=0; i < NumDevices; ++i) {
HIP_CHECK(hipSetDevice(i));
HipTest::initArrays(&A_d[i], &B_d[i], &C_d[i],
&A_h[i], &B_h[i], &C_h[i], N, false);
}
for (int i=0; i < NumDevices; ++i) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes,
hipMemcpyHostToDevice, stream[i]));
HIP_CHECK(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes,
hipMemcpyHostToDevice, stream[i]));
}
for (int i=0; i < NumDevices; ++i) {
HIP_CHECK(hipSetDevice(i));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks),
dim3(threadsPerBlock), 0, stream[i],
static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
}
for (int i=0; i < NumDevices; ++i) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipStreamSynchronize(stream[i]));
HIP_CHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
HipTest::checkVectorADD(A_h[i], B_h[i], C_h[i], N);
}
for (int i=0; i < NumDevices; ++i) {
HIP_CHECK(hipSetDevice(i));
HipTest::freeArrays(A_d[i], B_d[i], C_d[i],
A_h[i], B_h[i], C_h[i], false);
HIP_CHECK(hipStreamDestroy(stream[i]));
}
}
}
void TestkindDtoH(void) {
size_t Nbytes = N * sizeof(int);
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemcpyWithStream(A_d, A_h, Nbytes,
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipMemcpyWithStream(B_d, B_h, Nbytes,
hipMemcpyHostToDevice, stream));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipMemcpyWithStream(C_h, C_d, Nbytes,
hipMemcpyDeviceToHost, stream));
HipTest::checkVectorADD(A_h, B_h, C_h, N);
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipStreamDestroy(stream));
}
void TestkindDtoD(void) {
size_t Nbytes = N * sizeof(int);
int NumDevices = 0;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGetDeviceCount(&NumDevices));
// If you have single GPU machine the return
if (NumDevices <= 1) {
SUCCEED("NumDevices are less than 2");
} else {
int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices];
int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices];
hipStream_t stream[MaxGPUDevices];
for (int i=0; i < NumDevices; ++i) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipStreamCreate(&stream[i]));
}
// Initialize and create the host and device elements for first device
HIP_CHECK(hipSetDevice(0));
HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0],
&A_h[0], &B_h[0], &C_h[0], N, false);
for (int i=1; i < NumDevices; ++i) {
HIP_CHECK(hipSetDevice(i))
HIP_CHECK(hipMalloc(&A_d[i], Nbytes));
HIP_CHECK(hipMalloc(&B_d[i], Nbytes));
HIP_CHECK(hipMalloc(&C_d[i], Nbytes));
C_h[i] = reinterpret_cast<int*>(malloc(Nbytes));
HIP_ASSERT(C_h[i] != NULL);
}
HIP_CHECK(hipSetDevice(0));
HIP_CHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes,
hipMemcpyHostToDevice, stream[0]));
HIP_CHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes,
hipMemcpyHostToDevice, stream[0]));
// Copying device data from 1st GPU to the rest of the the GPUs that is
// NumDevices in the setup. 1st GPU start numbering from 0,1,2..n etc.
for (int i=1; i < NumDevices; ++i) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes,
hipMemcpyDeviceToDevice, stream[i]));
HIP_CHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes,
hipMemcpyDeviceToDevice, stream[i]));
}
// Launching the kernel including the 1st GPU to the no of GPUs present
// in the setup. 1st GPU start numbering from 0,1,2..n etc.
for (int i=0; i < NumDevices; ++i) {
HIP_CHECK(hipSetDevice(i));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks),
dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
}
for (int i=0; i < NumDevices; ++i) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipStreamSynchronize(stream[i]));
HIP_CHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N);
}
HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false);
HIP_CHECK(hipStreamDestroy(stream[0]));
for (int i=1; i < NumDevices; ++i) {
if (A_d[i]) {
HIP_CHECK(hipFree(A_d[i]));
}
if (B_d[i]) {
HIP_CHECK(hipFree(B_d[i]));
}
if (C_d[i]) {
HIP_CHECK(hipFree(C_d[i]));
}
if (C_h[i]) {
free(C_h[i]);
}
HIP_CHECK(hipStreamDestroy(stream[i]));
}
}
}
void TestkindDefault(void) {
size_t Nbytes = N * sizeof(int);
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyDefault, stream));
HIP_CHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyDefault, stream));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDefault, stream));
HipTest::checkVectorADD(A_h, B_h, C_h, N);
HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
HIP_CHECK(hipStreamDestroy(stream));
}
void TestkindDefaultForDtoD(void) {
size_t Nbytes = N * sizeof(int);
int NumDevices = 0;
unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIP_CHECK(hipGetDeviceCount(&NumDevices));
// Test case will not run on single GPU setup.
if (NumDevices <= 1) {
SUCCEED("No of Devices < 2");
} else {
int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices];
int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices];
// Initialize and create the host and device elements for first device
HIP_CHECK(hipSetDevice(0));
HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0],
&A_h[0], &B_h[0], &C_h[0], N, false);
for (int i=1; i < NumDevices; ++i) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipMalloc(&A_d[i], Nbytes));
HIP_CHECK(hipMalloc(&B_d[i], Nbytes));
HIP_CHECK(hipMalloc(&C_d[i], Nbytes));
C_h[i] = reinterpret_cast<int*>(malloc(Nbytes));
HIP_ASSERT(C_h[i] != NULL);
}
hipStream_t stream[MaxGPUDevices];
for (int i=0; i < NumDevices; ++i) {
HIP_CHECK(hipSetDevice(i));
HIP_CHECK(hipStreamCreate(&stream[i]));
}
HIP_CHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes,
hipMemcpyHostToDevice, stream[0]));
HIP_CHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes,
hipMemcpyHostToDevice, stream[0]));
// Copying device data from 1st GPU to the rest of the the GPUs
// using hipMemcpyDefault kind that is NumDevices in the setup.
// 1st GPU start numbering from 0,1,2..n etc.
for (int i=1; i < NumDevices; ++i) {
HIP_CHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes,
hipMemcpyDefault, stream[i]));
HIP_CHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes,
hipMemcpyDefault, stream[i]));
}
for (int i=0; i < NumDevices; ++i) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks),
dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
}
for (int i=0; i < NumDevices; ++i) {
HIP_CHECK(hipSetDevice(i)); // hipMemcpy will be on this device
HIP_CHECK(hipStreamSynchronize(stream[i]));
HIP_CHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
// Output of each GPU is getting validated with input of 1st GPU.
HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N);
}
HipTest::freeArrays(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0], false);
HIP_CHECK(hipStreamDestroy(stream[0]));
for (int i=1; i < NumDevices; ++i) {
if (A_d[i]) {
HIP_CHECK(hipFree(A_d[i]));
}
if (B_d[i]) {
HIP_CHECK(hipFree(B_d[i]));
}
if (C_d[i]) {
HIP_CHECK(hipFree(C_d[i]));
}
if (C_h[i]) {
free(C_h[i]);
}
HIP_CHECK(hipStreamDestroy(stream[i]));
}
}
}
void TestkindHtoH(void) {
size_t Nbytes = N * sizeof(int);
int *A_h, *B_h;
// Allocate memory to A_h and B_h
A_h = static_cast<int*>(malloc(Nbytes));
HIP_ASSERT(A_h != NULL);
B_h = static_cast<int*>(malloc(Nbytes));
HIP_ASSERT(B_h != NULL);
for (size_t i = 0; i < N; ++i) {
if (A_h) {
(A_h)[i] = 3.146f + i; // Pi
}
}
hipStream_t stream;
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(hipMemcpyWithStream(B_h, A_h, Nbytes, hipMemcpyHostToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
for (size_t i = 0; i < N; i++) {
HIP_ASSERT(A_h[i] == B_h[i]);
}
if (A_h) {
free(A_h);
}
if (B_h) {
free(B_h);
}
HIP_CHECK(hipStreamDestroy(stream));
}
TEST_CASE("Unit_hipMemcpyWithStream_MultiThread") {
size_t thread_count = 10;
std::vector<joinable_thread> threads;
int deviceCount = 0;
HIP_CHECK(hipGetDeviceCount(&deviceCount));
if (deviceCount < 2) {
SUCCEED("deviceCount < 2");
} else {
for (int op = static_cast<int>(ops::TestwithOnestream);
op < static_cast<int>(ops::END_OF_LIST); ++op) {
for (uint32_t i = 0; i < thread_count; i++) {
threads.emplace_back(std::thread{[&] {
switch ( op ) {
case static_cast<int>(ops::TestwithOnestream):
TestwithOnestream();
break;
case static_cast<int>(ops::TestwithTwoStream):
TestwithTwoStream();
break;
case static_cast<int>(ops::TestkindDtoH):
TestkindDtoH();
break;
case static_cast<int>(ops::TestkindHtoH):
TestkindHtoH();
break;
case static_cast<int>(ops::TestkindDtoD):
TestkindDtoD();
break;
case static_cast<int>(ops::TestOnMultiGPUwithOneStream):
TestOnMultiGPUwithOneStream();
break;
case static_cast<int>(ops::TestkindDefault):
TestkindDefault();
break;
#ifndef __HIP_PLATFORM_NVCC__
case static_cast<int>(ops::TestkindDefaultForDtoD):
TestkindDefaultForDtoD();
break;
#endif
case static_cast<int>(ops::TestDtoDonSameDevice):
TestDtoDonSameDevice();
break;
default:{}
}
}});
}
}
}
}
TEST_CASE("Unit_hipMemcpyWithStream_TestWithOneStream") {
TestwithOnestream();
}
TEST_CASE("Unit_hipMemcpyWithStream_TestwithTwoStream") {
TestwithTwoStream();
}
TEST_CASE("Unit_hipMemcpyWithStream_TestkindDtoH") {
TestkindDtoH();
}
TEST_CASE("Unit_hipMemcpyWithStream_TestkindHtoH") {
TestkindHtoH();
}
TEST_CASE("Unit_hipMemcpyWithStream_TestkindDtoD") {
TestkindDtoD();
}
TEST_CASE("Unit_hipMemcpyWithStream_TestOnMultiGPUwithOneStream") {
TestOnMultiGPUwithOneStream();
}
TEST_CASE("Unit_hipMemcpyWithStream_TestkindDefault") {
TestkindDefault();
}
#ifndef __HIP_PLATFORM_NVCC__
TEST_CASE("Unit_hipMemcpyWithStream_TestkindDefaultForDtoD") {
TestkindDefaultForDtoD();
}
#endif
TEST_CASE("Unit_hipMemcpyWithStream_TestDtoDonSameDevice") {
TestDtoDonSameDevice();
}
catch/unit/memory/hipMemcpyWithStreamMultiThread.cc
+679
View File
@@ -0,0 +1,679 @@
/*
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.
*/
/*
* Different test for checking functionality of
* hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,
* hipMemcpyKind kind, hipStream_t stream);
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <vector>
#define LEN 64
#define SIZE LEN << 2
#define THREADS 2
#define MAX_THREADS 16
static constexpr auto N{1024};
static constexpr auto Nbytes{N * sizeof(int)};
static const auto MaxGPUDevices{256};
static constexpr unsigned blocksPerCU{6}; // to hide latency
static constexpr unsigned threadsPerBlock{256};
enum class ops
{ TestwithOnestream,
TestwithTwoStream,
TestOnMultiGPUwithOneStream,
TestkindDtoH,
TestkindDtoD,
TestkindHtoH,
TestkindDefault,
TestkindDefaultForDtoD,
TestDtoDonSameDevice,
END_OF_LIST
};
namespace MemcpyStream {
unsigned setNumBlocks(int blocksPerCU, int 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 MemcpyStream
class HipMemcpyWithStreamMultiThreadtests {
public:
// Test hipMemcpyWithStream with one streams and launch kernel in
// that stream, verify the data.
void TestwithOnestream(bool &val_res);
// Test hipMemcpyWithStream with two streams and launch kernels in
// two streams, verify the data.
void TestwithTwoStream(bool &val_res);
// Test hipMemcpyWithStream with one stream for each gpu and launch
// kernels in each, verify the data
void TestOnMultiGPUwithOneStream(bool &val_res);
// Test hipMemcpyWithStream to copy data from
// device to host (hipMemcpyDeviceToHost).
void TestkindDtoH(bool &val_res);
// Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on MultiGPU.
void TestkindDtoD(bool &val_res);
// Test hipMemcpyWithStream with hipMemcpyHostToHost.
void TestkindHtoH(bool &val_res);
// Test hipMemcpyWithStream with hipMemcpyDefault.
void TestkindDefault(bool &val_res);
// Test hipMemcpyWithStream with hipMemcpyDefault for
// device to device transfer case.
void TestkindDefaultForDtoD(bool &val_res);
// Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on same device.
void TestDtoDonSameDevice(bool &val_res);
// Allocate Memory
void AllocateMemory(int** A_d, int** B_d,
int** C_d, int** A_h,
int** B_h,
int** C_h);
// DeAllocate Memory
void DeAllocateMemory(int* A_d, int* B_d,
int* C_d, int* A_h, int* B_h,
int* C_h);
// Validate Result
bool ValidateResult(int *A_h, int *B_h, int *C_h);
};
void HipMemcpyWithStreamMultiThreadtests::AllocateMemory(int** A_d, int** B_d,
int** C_d, int** A_h,
int** B_h,
int** C_h) {
HIPCHECK(hipMalloc(A_d, Nbytes));
HIPCHECK(hipMalloc(B_d, Nbytes));
HIPCHECK(hipMalloc(C_d, Nbytes));
*A_h = reinterpret_cast<int*>(malloc(Nbytes));
*B_h = reinterpret_cast<int*>(malloc(Nbytes));
*C_h = reinterpret_cast<int*>(malloc(Nbytes));
for (size_t i = 0; i < N; i++) {
if (*A_h) (*A_h)[i] = 3;
if (*B_h) (*B_h)[i] = 4;
if (*C_h) (*C_h)[i] = 5;
}
}
void HipMemcpyWithStreamMultiThreadtests::DeAllocateMemory(int* A_d, int* B_d,
int* C_d, int* A_h, int* B_h,
int* C_h) {
HIP_CHECK(hipFree(A_d));
HIP_CHECK(hipFree(B_d));
HIP_CHECK(hipFree(C_d));
free(A_h);
free(B_h);
free(C_h);
}
bool HipMemcpyWithStreamMultiThreadtests::
ValidateResult(int *A_h, int *B_h, int *C_h) {
bool TestPassed = true;
for (size_t i = 0; i < N; i++) {
if ((A_h[i] + B_h[i]) != C_h[i]) {
TestPassed = false;
break;
}
}
return TestPassed;
}
void HipMemcpyWithStreamMultiThreadtests::TestwithOnestream(bool &val_res) {
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
size_t Nbytes{N * sizeof(int)};
AllocateMemory(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h);
unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N);
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes,
hipMemcpyHostToDevice, stream));
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes,
hipMemcpyHostToDevice, stream));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
HIPCHECK(hipStreamSynchronize(stream));
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
val_res = ValidateResult(A_h, B_h, C_h);
DeAllocateMemory(A_d, B_d, C_d, A_h, B_h, C_h);
HIPCHECK(hipStreamDestroy(stream));
}
void HipMemcpyWithStreamMultiThreadtests::TestwithTwoStream(bool &val_res) {
size_t Nbytes = N * sizeof(int);
const int NoofStreams = 2;
int *A_d[NoofStreams], *B_d[NoofStreams], *C_d[NoofStreams];
int *A_h[NoofStreams], *B_h[NoofStreams], *C_h[NoofStreams];
unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N);
for (int i=0; i < NoofStreams; ++i) {
AllocateMemory(&A_d[i], &B_d[i], &C_d[i],
&A_h[i], &B_h[i], &C_h[i]);
}
hipStream_t stream[NoofStreams];
for (int i=0; i < NoofStreams; ++i) {
HIPCHECK(hipStreamCreate(&stream[i]));
}
for (int i=0; i < NoofStreams; ++i) {
HIPCHECK(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes,
hipMemcpyHostToDevice, stream[i]));
HIPCHECK(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes,
hipMemcpyHostToDevice, stream[i]));
}
for (int i=0; i < NoofStreams; ++i) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
}
for (int i=0; i < NoofStreams; ++i) {
HIPCHECK(hipStreamSynchronize(stream[i]));
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
val_res = ValidateResult(A_h[i], B_h[i], C_h[i]);
}
for (int i=0; i < NoofStreams; ++i) {
DeAllocateMemory(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i]);
HIPCHECK(hipStreamDestroy(stream[i]));
}
}
void HipMemcpyWithStreamMultiThreadtests::TestDtoDonSameDevice(bool &val_res) {
size_t Nbytes = N * sizeof(int);
const int NoofStreams = 2;
int *A_d[NoofStreams], *B_d[NoofStreams], *C_d[NoofStreams];
int *A_h[NoofStreams], *B_h[NoofStreams], *C_h[NoofStreams];
unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N);
AllocateMemory(&A_d[0], &B_d[0], &C_d[0],
&A_h[0], &B_h[0], &C_h[0]);
hipStream_t stream[NoofStreams];
for (int i=0; i < NoofStreams; ++i) {
HIPCHECK(hipSetDevice(0));
HIPCHECK(hipStreamCreate(&stream[i]));
}
HIPCHECK(hipSetDevice(0));
HIPCHECK(hipMalloc(&A_d[1], Nbytes));
HIPCHECK(hipMalloc(&B_d[1], Nbytes));
HIPCHECK(hipMalloc(&C_d[1], Nbytes));
C_h[1] = reinterpret_cast<int*>(malloc(Nbytes));
HIPASSERT(C_h[1] != NULL);
HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes,
hipMemcpyHostToDevice, stream[0]));
HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes,
hipMemcpyHostToDevice, stream[0]));
HIPCHECK(hipMemcpyWithStream(A_d[1], A_d[0], Nbytes,
hipMemcpyDeviceToDevice, stream[1]));
HIPCHECK(hipMemcpyWithStream(B_d[1], B_d[0], Nbytes,
hipMemcpyDeviceToDevice, stream[1]));
for (int i=0; i < NoofStreams; ++i) {
HIPCHECK(hipSetDevice(0));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
}
for (int i=0; i < NoofStreams; ++i) {
HIPCHECK(hipSetDevice(0));
HIPCHECK(hipStreamSynchronize(stream[i]));
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
val_res = ValidateResult(A_h[0], B_h[0], C_h[i]);
}
DeAllocateMemory(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0]);
if (A_d[1]) {
HIPCHECK(hipFree(A_d[1]));
}
if (B_d[1]) {
HIPCHECK(hipFree(B_d[1]));
}
if (C_d[1]) {
HIPCHECK(hipFree(C_d[1]));
}
if (C_h[1]) {
free(C_h[1]);
}
for (int i=0; i < NoofStreams; ++i) {
HIPCHECK(hipStreamDestroy(stream[i]));
}
}
void HipMemcpyWithStreamMultiThreadtests::
TestOnMultiGPUwithOneStream(bool &val_res) {
size_t Nbytes = N * sizeof(int);
int numDevices = 0;
unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIPCHECK(hipGetDeviceCount(&numDevices));
// If you have single GPU machine the return
if (numDevices <= 1) {
return;
}
int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices];
int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices];
hipStream_t stream[MaxGPUDevices];
for (int i=0; i < numDevices; ++i) {
HIPCHECK(hipSetDevice(i));
HIPCHECK(hipStreamCreate(&stream[i]));
}
for (int i=0; i < numDevices; ++i) {
HIPCHECK(hipSetDevice(i));
AllocateMemory(&A_d[i], &B_d[i], &C_d[i],
&A_h[i], &B_h[i], &C_h[i]);
}
for (int i=0; i < numDevices; ++i) {
HIPCHECK(hipSetDevice(i));
HIPCHECK(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes,
hipMemcpyHostToDevice, stream[i]));
HIPCHECK(hipMemcpyWithStream(B_d[i], B_h[i], Nbytes,
hipMemcpyHostToDevice, stream[i]));
}
for (int i=0; i < numDevices; ++i) {
HIPCHECK(hipSetDevice(i));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
}
for (int i=0; i < numDevices; ++i) {
HIPCHECK(hipSetDevice(i));
HIPCHECK(hipStreamSynchronize(stream[i]));
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
val_res = ValidateResult(A_h[i], B_h[i], C_h[i]);
}
for (int i=0; i < numDevices; ++i) {
HIPCHECK(hipSetDevice(i));
DeAllocateMemory(A_d[i], B_d[i], C_d[i], A_h[i], B_h[i], C_h[i]);
HIPCHECK(hipStreamDestroy(stream[i]));
}
}
void HipMemcpyWithStreamMultiThreadtests::TestkindDtoH(bool &val_res) {
size_t Nbytes = N * sizeof(int);
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N);
AllocateMemory(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h);
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes,
hipMemcpyHostToDevice, stream));
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes,
hipMemcpyHostToDevice, stream));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
HIPCHECK(hipStreamSynchronize(stream));
HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes,
hipMemcpyDeviceToHost, stream));
val_res = ValidateResult(A_h, B_h, C_h);
DeAllocateMemory(A_d, B_d, C_d, A_h, B_h, C_h);
HIPCHECK(hipStreamDestroy(stream));
}
void HipMemcpyWithStreamMultiThreadtests::TestkindDtoD(bool &val_res) {
size_t Nbytes = N * sizeof(int);
int numDevices = 0;
unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIPCHECK(hipGetDeviceCount(&numDevices));
// If you have single GPU machine the return
if (numDevices <= 1) {
return;
}
int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices];
int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices];
hipStream_t stream[MaxGPUDevices];
for (int i=0; i < numDevices; ++i) {
HIPCHECK(hipSetDevice(i));
HIPCHECK(hipStreamCreate(&stream[i]));
}
// Initialize and create the host and device elements for first device
HIPCHECK(hipSetDevice(0));
AllocateMemory(&A_d[0], &B_d[0], &C_d[0],
&A_h[0], &B_h[0], &C_h[0]);
for (int i=1; i < numDevices; ++i) {
HIPCHECK(hipSetDevice(i))
HIPCHECK(hipMalloc(&A_d[i], Nbytes));
HIPCHECK(hipMalloc(&B_d[i], Nbytes));
HIPCHECK(hipMalloc(&C_d[i], Nbytes));
C_h[i] = reinterpret_cast<int*>(malloc(Nbytes));
HIPASSERT(C_h[i] != NULL);
}
HIPCHECK(hipSetDevice(0));
HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes,
hipMemcpyHostToDevice, stream[0]));
HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes,
hipMemcpyHostToDevice, stream[0]));
// Copying device data from 1st GPU to the rest of the the GPUs that is
// numDevices in the setup. 1st GPU start numbering from 0,1,2..n etc.
for (int i=1; i < numDevices; ++i) {
HIPCHECK(hipSetDevice(i));
HIPCHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes,
hipMemcpyDeviceToDevice, stream[i]));
HIPCHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes,
hipMemcpyDeviceToDevice, stream[i]));
}
// Launching the kernel including the 1st GPU to the no of GPUs present
// in the setup. 1st GPU start numbering from 0,1,2..n etc.
for (int i=0; i < numDevices; ++i) {
HIPCHECK(hipSetDevice(i));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
}
for (int i=0; i < numDevices; ++i) {
HIPCHECK(hipSetDevice(i));
HIPCHECK(hipStreamSynchronize(stream[i]));
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
val_res = ValidateResult(A_h[0], B_h[0], C_h[i]);
}
DeAllocateMemory(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0]);
HIPCHECK(hipStreamDestroy(stream[0]));
for (int i=1; i < numDevices; ++i) {
if (A_d[i]) {
HIPCHECK(hipFree(A_d[i]));
}
if (B_d[i]) {
HIPCHECK(hipFree(B_d[i]));
}
if (C_d[i]) {
HIPCHECK(hipFree(C_d[i]));
}
if (C_h[i]) {
free(C_h[i]);
}
HIPCHECK(hipStreamDestroy(stream[i]));
}
}
void HipMemcpyWithStreamMultiThreadtests::
TestkindDefault(bool &val_res) {
size_t Nbytes = N * sizeof(int);
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N);
AllocateMemory(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h);
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
HIPCHECK(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyDefault, stream));
HIPCHECK(hipMemcpyWithStream(B_d, B_h, Nbytes, hipMemcpyDefault, stream));
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream, static_cast<const int*>(A_d),
static_cast<const int*>(B_d), C_d, N);
HIPCHECK(hipStreamSynchronize(stream));
HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDefault, stream));
val_res = ValidateResult(A_h, B_h, C_h);
DeAllocateMemory(A_d, B_d, C_d, A_h, B_h, C_h);
HIPCHECK(hipStreamDestroy(stream));
}
void HipMemcpyWithStreamMultiThreadtests::
TestkindDefaultForDtoD(bool &val_res) {
size_t Nbytes = N * sizeof(int);
int numDevices = 0;
unsigned blocks = MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N);
HIPCHECK(hipGetDeviceCount(&numDevices));
// Test case will not run on single GPU setup.
if (numDevices <= 1) {
return;
}
int *A_d[MaxGPUDevices], *B_d[MaxGPUDevices], *C_d[MaxGPUDevices];
int *A_h[MaxGPUDevices], *B_h[MaxGPUDevices], *C_h[MaxGPUDevices];
// Initialize and create the host and device elements for first device
HIPCHECK(hipSetDevice(0));
AllocateMemory(&A_d[0], &B_d[0], &C_d[0],
&A_h[0], &B_h[0], &C_h[0]);
for (int i=1; i < numDevices; ++i) {
HIPCHECK(hipSetDevice(i));
HIPCHECK(hipMalloc(&A_d[i], Nbytes));
HIPCHECK(hipMalloc(&B_d[i], Nbytes));
HIPCHECK(hipMalloc(&C_d[i], Nbytes));
C_h[i] = reinterpret_cast<int*>(malloc(Nbytes));
HIPASSERT(C_h[i] != NULL);
}
hipStream_t stream[MaxGPUDevices];
for (int i=0; i < numDevices; ++i) {
HIPCHECK(hipSetDevice(i));
HIPCHECK(hipStreamCreate(&stream[i]));
}
HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes,
hipMemcpyHostToDevice, stream[0]));
HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes,
hipMemcpyHostToDevice, stream[0]));
// Copying device data from 1st GPU to the rest of the the GPUs
// using hipMemcpyDefault kind that is numDevices in the setup.
// 1st GPU start numbering from 0,1,2..n etc.
for (int i=1; i < numDevices; ++i) {
HIPCHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes,
hipMemcpyDefault, stream[i]));
HIPCHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes,
hipMemcpyDefault, stream[i]));
}
for (int i=0; i < numDevices; ++i) {
hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock),
0, stream[i], static_cast<const int*>(A_d[i]),
static_cast<const int*>(B_d[i]), C_d[i], N);
}
for (int i=0; i < numDevices; ++i) {
HIPCHECK(hipSetDevice(i)); // hipMemcpy will be on this device
HIPCHECK(hipStreamSynchronize(stream[i]));
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
// Output of each GPU is getting validated with input of 1st GPU.
val_res = ValidateResult(A_h[0], B_h[0], C_h[i]);
}
DeAllocateMemory(A_d[0], B_d[0], C_d[0], A_h[0], B_h[0], C_h[0]);
HIPCHECK(hipStreamDestroy(stream[0]));
for (int i=1; i < numDevices; ++i) {
if (A_d[i]) {
HIPCHECK(hipFree(A_d[i]));
}
if (B_d[i]) {
HIPCHECK(hipFree(B_d[i]));
}
if (C_d[i]) {
HIPCHECK(hipFree(C_d[i]));
}
if (C_h[i]) {
free(C_h[i]);
}
HIPCHECK(hipStreamDestroy(stream[i]));
}
}
void HipMemcpyWithStreamMultiThreadtests::TestkindHtoH(bool &val_res) {
size_t Nbytes = N * sizeof(int);
int *A_h, *B_h;
// Allocate memory to A_h and B_h
A_h = static_cast<int*>(malloc(Nbytes));
B_h = static_cast<int*>(malloc(Nbytes));
for (size_t i = 0; i < N; ++i) {
if (A_h) {
(A_h)[i] = 3.146f + i; // Pi
}
}
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
HIPCHECK(hipMemcpyWithStream(B_h, A_h, Nbytes, hipMemcpyHostToHost, stream));
HIPCHECK(hipStreamSynchronize(stream));
for (size_t i = 0; i < N; i++) {
if ((A_h[i] != B_h[i])) {
val_res = false;
break;
}
}
if (A_h) {
free(A_h);
}
if (B_h) {
free(B_h);
}
HIPCHECK(hipStreamDestroy(stream));
}
TEST_CASE("Unit_hipMemcpyWithStream_NewMultiThread") {
const auto Threadcount{100};
bool ret_val[Threadcount];
std::thread th[Threadcount];
for (int op = static_cast<int>(ops::TestwithOnestream);
op < static_cast<int>(ops::END_OF_LIST); ++op) {
HipMemcpyWithStreamMultiThreadtests tests;
for (uint32_t i = 0; i < Threadcount; i++) {
switch ( static_cast<ops>(op) ) {
case ops::TestwithOnestream:
th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::
TestwithOnestream,
&tests, std::ref(ret_val[i]));
break;
case ops::TestwithTwoStream:
th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::
TestwithTwoStream,
&tests, std::ref(ret_val[i]));
break;
case ops::TestkindDtoH:
th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::
TestkindDtoH,
&tests, std::ref(ret_val[i]));
break;
case ops::TestkindHtoH:
th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::
TestkindHtoH,
&tests, std::ref(ret_val[i]));
break;
case ops::TestkindDtoD:
th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::
TestkindDtoD,
&tests, std::ref(ret_val[i]));
break;
case ops::TestOnMultiGPUwithOneStream:
th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::
TestOnMultiGPUwithOneStream,
&tests, std::ref(ret_val[i]));
break;
case ops::TestkindDefault:
th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::
TestkindDefault,
&tests, std::ref(ret_val[i]));
break;
case ops::TestkindDefaultForDtoD:
th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::
TestkindDefaultForDtoD,
&tests, std::ref(ret_val[i]));
break;
case ops::TestDtoDonSameDevice:
th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::
TestDtoDonSameDevice,
&tests, std::ref(ret_val[i]));
break;
default: {}
}
}
for (uint32_t i = 0; i < Threadcount; i++) {
th[i].join();
}
for (uint32_t i = 0; i < Threadcount; i++) {
REQUIRE(ret_val[i] == true);
}
}
}
catch/unit/memory/hipMemcpy_MultiThread.cc
+322
View File
@@ -0,0 +1,322 @@
/*
Copyright (c) 2021 - 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.
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
#include <utility>
#include <vector>
/*
This testfile verifies the following scenarios of all hipMemcpy API
1. Multi thread
*/
static constexpr auto NUM_ELM{1024};
static constexpr auto NUM_THREADS{10};
static auto Available_Gpus{0};
static constexpr auto MAX_GPU{256};
enum apiToTest {TEST_MEMCPY, TEST_MEMCPYH2D, TEST_MEMCPYD2H, TEST_MEMCPYD2D,
TEST_MEMCPYASYNC, TEST_MEMCPYH2DASYNC, TEST_MEMCPYD2HASYNC,
TEST_MEMCPYD2DASYNC};
template <typename T>
class memcpyTests {
public:
T *A_h, *B_h;
apiToTest api;
explicit memcpyTests(apiToTest val);
memcpyTests() = delete;
void Memcpy_And_verify(bool *ret_val);
bool CheckTests(T* A_h, T* B_h, int NUM_ELEMENTS);
~memcpyTests();
};
template <typename T>
bool memcpyTests<T>::CheckTests(T *A_h, T *B_h, int NUM_ELEMENTS) {
for (auto i =0; i < NUM_ELEMENTS; i++) {
if (A_h[i] != B_h[i]) {
return false;
}
}
return true;
}
template <typename T>
memcpyTests<T>::memcpyTests(apiToTest val) {
api = val;
A_h = reinterpret_cast<T*>(malloc(NUM_ELM * sizeof(T)));
B_h = reinterpret_cast<T*>(malloc(NUM_ELM * sizeof(T)));
if ((A_h == nullptr) || (B_h == nullptr)) {
exit(1);
}
for (size_t i = 0; i < NUM_ELM; ++i) {
A_h[i] = 123;
B_h[i] = 0;
}
}
template <typename T>
void memcpyTests<T>::Memcpy_And_verify(bool *ret_val) {
HIPCHECK(hipGetDeviceCount(&Available_Gpus));
T *A_d[MAX_GPU];
hipStream_t stream[MAX_GPU];
for (int i = 0; i < Available_Gpus; ++i) {
HIPCHECK(hipSetDevice(i));
HIPCHECK(hipMalloc(&A_d[i], NUM_ELM * sizeof(T)));
if (api >= TEST_MEMCPYD2D) {
HIPCHECK(hipStreamCreate(&stream[i]));
}
}
HIPCHECK(hipSetDevice(0));
int canAccessPeer = 0;
switch (api) {
case TEST_MEMCPY:
{
// To test hipMemcpy()
// Copying data from host to individual devices followed by copying
// back to host and verifying the data consistency.
for (int i = 0; i < Available_Gpus; ++i) {
HIPCHECK(hipMemcpy(A_d[i], A_h, NUM_ELM * sizeof(T),
hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(T),
hipMemcpyDeviceToHost));
*ret_val = CheckTests(A_h, B_h, NUM_ELM);
}
// Device to Device copying for all combinations
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
if (canAccessPeer) {
HIPCHECK(hipMemcpy(A_d[j], A_d[i], NUM_ELM * sizeof(T),
hipMemcpyDefault));
// Copying in reverse dir of above to check if bidirectional
// access is happening without any error
HIPCHECK(hipMemcpy(A_d[i], A_d[j], NUM_ELM * sizeof(T),
hipMemcpyDefault));
// Copying data to host to verify the content
HIPCHECK(hipMemcpy(B_h, A_d[j], NUM_ELM * sizeof(T),
hipMemcpyDefault));
*ret_val &= CheckTests(A_h, B_h, NUM_ELM);
}
}
}
break;
}
case TEST_MEMCPYH2D: // To test hipMemcpyHtoD()
{
for (int i = 0; i < Available_Gpus; ++i) {
HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[i]),
A_h, NUM_ELM * sizeof(T)));
// Copying data from device to host to check data consistency
HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(T),
hipMemcpyDeviceToHost));
*ret_val &= CheckTests(A_h, B_h, NUM_ELM);
}
break;
}
case TEST_MEMCPYD2H: // To test hipMemcpyDtoH()--done
{
for (int i = 0; i < Available_Gpus; ++i) {
HIPCHECK(hipMemcpy(A_d[i], A_h, NUM_ELM * sizeof(T),
hipMemcpyHostToDevice));
HIPCHECK(hipMemcpyDtoH(B_h, hipDeviceptr_t(A_d[i]),
NUM_ELM * sizeof(T)));
*ret_val &= CheckTests(A_h, B_h, NUM_ELM);
}
break;
}
case TEST_MEMCPYD2D: // To test hipMemcpyDtoD()
{
if (Available_Gpus > 1) {
// First copy data from H to D and then
// from D to D followed by D to H
// HIPCHECK(hipMemcpyHtoD(A_d[0], A_h,
// NUM_ELM * sizeof(T)));
int canAccessPeer = 0;
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
if (canAccessPeer) {
HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[i]),
A_h, NUM_ELM * sizeof(T)));
HIPCHECK(hipMemcpyDtoD(hipDeviceptr_t(A_d[j]),
hipDeviceptr_t(A_d[i]), NUM_ELM * sizeof(T)));
// Copying in direction reverse of above to check if
// bidirectional
// access is happening without any error
HIPCHECK(hipMemcpyDtoD(hipDeviceptr_t(A_d[i]),
hipDeviceptr_t(A_d[j]), NUM_ELM * sizeof(T)));
HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(T),
hipMemcpyDeviceToHost));
*ret_val &= CheckTests(A_h, B_h, NUM_ELM);
}
}
}
} else {
// As DtoD is not possible transfer data from HtH(A_h to B_h)
// so as to get through verification step
HIPCHECK(hipMemcpy(B_h, A_h, NUM_ELM * sizeof(T),
hipMemcpyHostToHost));
*ret_val &= CheckTests(A_h, B_h, NUM_ELM);
}
break;
}
case TEST_MEMCPYASYNC:
{
// To test hipMemcpyAsync()
// Copying data from host to individual devices followed by copying
// back to host and verifying the data consistency.
for (int i = 0; i < Available_Gpus; ++i) {
HIPCHECK(hipMemcpyAsync(A_d[i], A_h, NUM_ELM * sizeof(T),
hipMemcpyHostToDevice, stream[i]));
HIPCHECK(hipMemcpyAsync(B_h, A_d[i], NUM_ELM * sizeof(T),
hipMemcpyDeviceToHost, stream[i]));
HIPCHECK(hipStreamSynchronize(stream[i]));
*ret_val &= CheckTests(A_h, B_h, NUM_ELM);
}
// Device to Device copying for all combinations
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
if (canAccessPeer) {
HIPCHECK(hipMemcpyAsync(A_d[j], A_d[i],
NUM_ELM * sizeof(T),
hipMemcpyDefault, stream[i]));
// Copying in direction reverse of above to
// check if bidirectional
// access is happening without any error
HIPCHECK(hipMemcpyAsync(A_d[i], A_d[j],
NUM_ELM * sizeof(T),
hipMemcpyDefault, stream[i]));
HIPCHECK(hipStreamSynchronize(stream[i]));
HIPCHECK(hipMemcpy(B_h, A_d[j], NUM_ELM * sizeof(T),
hipMemcpyDefault));
*ret_val &= CheckTests(A_h, B_h, NUM_ELM);
}
}
}
break;
}
case TEST_MEMCPYH2DASYNC: // To test hipMemcpyHtoDAsync()
{
for (int i = 0; i < Available_Gpus; ++i) {
HIPCHECK(hipMemcpyHtoDAsync(hipDeviceptr_t(A_d[i]), A_h,
NUM_ELM * sizeof(T), stream[i]));
HIPCHECK(hipStreamSynchronize(stream[i]));
// Copying data from device to host to check data consistency
HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(T),
hipMemcpyDeviceToHost));
*ret_val &= CheckTests(A_h, B_h, NUM_ELM);
}
break;
}
case TEST_MEMCPYD2HASYNC: // To test hipMemcpyDtoHAsync()
{
for (int i = 0; i < Available_Gpus; ++i) {
HIPCHECK(hipMemcpy(A_d[i], A_h, NUM_ELM * sizeof(T),
hipMemcpyHostToDevice));
HIPCHECK(hipMemcpyDtoHAsync(B_h, hipDeviceptr_t(A_d[i]),
NUM_ELM * sizeof(T), stream[i]));
HIPCHECK(hipStreamSynchronize(stream[i]));
*ret_val &= CheckTests(A_h, B_h, NUM_ELM);
}
break;
}
case TEST_MEMCPYD2DASYNC: // To test hipMemcpyDtoDAsync()
{
if (Available_Gpus > 1) {
// First copy data from H to D and then from D to D followed by D2H
HIPCHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[0]),
A_h, NUM_ELM * sizeof(T)));
for (int i = 0; i < Available_Gpus; ++i) {
for (int j = i+1; j < Available_Gpus; ++j) {
canAccessPeer = 0;
hipDeviceCanAccessPeer(&canAccessPeer, i, j);
if (canAccessPeer) {
HIPCHECK(hipSetDevice(j));
HIPCHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d[j]),
hipDeviceptr_t(A_d[i]), NUM_ELM * sizeof(T),
stream[i]));
// Copying in direction reverse of above to check if
// bidirectional
// access is happening without any error
HIPCHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d[i]),
hipDeviceptr_t(A_d[j]), NUM_ELM * sizeof(T),
stream[i]));
HIPCHECK(hipStreamSynchronize(stream[i]));
HIPCHECK(hipMemcpy(B_h, A_d[i], NUM_ELM * sizeof(T),
hipMemcpyDeviceToHost));
*ret_val &= CheckTests(A_h, B_h, NUM_ELM);
}
}
}
} else {
// As DtoD is not possible we will transfer data
// from HtH(A_h to B_h)
// so as to get through verification step
HIPCHECK(hipMemcpy(B_h, A_h, NUM_ELM * sizeof(T),
hipMemcpyHostToHost));
*ret_val &= CheckTests(A_h, B_h, NUM_ELM);
}
break;
}
}
for (int i = 0; i < Available_Gpus; ++i) {
HIPCHECK(hipSetDevice(i));
HIPCHECK(hipFree((A_d[i])));
if (api >= TEST_MEMCPYD2D) {
HIPCHECK(hipStreamDestroy(stream[i]));
}
}
}
template <typename T>
memcpyTests<T>::~memcpyTests() {
free(A_h);
free(B_h);
}
void Thread_func(bool &ret_val) {
for (apiToTest api = TEST_MEMCPY; api <= TEST_MEMCPYD2DASYNC;
api = apiToTest(api + 1)) {
memcpyTests<int> obj(api);
obj.Memcpy_And_verify(&ret_val);
}
}
TEST_CASE("Unit_hipMemcpy_MultiThread-AllAPIs") {
std::thread Thrd[NUM_THREADS];
bool ret_val[NUM_THREADS];
for (int i = 0; i < NUM_THREADS; i++)
Thrd[i] = std::thread(Thread_func, std::ref(ret_val[i]));
// Thread join is being called separately so as to allow the
// threads run parallely
for (int i = 0; i < NUM_THREADS; i++)
Thrd[i].join();
for (int i = 0; i < NUM_THREADS; i++)
REQUIRE(ret_val[i] == true);
}