SWDEV-398975 - re-enable Unit_hipMemcpyWithStream_MultiThread test
use correct macro inside the test
Change-Id: I7034a97efe91d6c9e2d9b4e6026ec770714a01ac
[ROCm/hip-tests commit: 9704c269d0]
This commit is contained in:
Jatin Chaudhary
committad av
Jatin Jaikishan Chaudhary
Jatin Jaikishan Chaudhary
förälder
e473199c50
incheckning
f09b9af969
@@ -92,8 +92,6 @@
|
||||
"Unit_hipMemcpyPeer_Positive_ZeroSize",
|
||||
"Unit_hipMemcpyPeerAsync_Positive_ZeroSize",
|
||||
"Fails in Stress test SWDEV-398971",
|
||||
"SWDEV-398975 Seg faults in stress test",
|
||||
"Unit_hipMemcpyWithStream_MultiThread",
|
||||
"SWDEV-398977 fails in stress tests",
|
||||
"Unit_hipMemset2DSync",
|
||||
"SWDEV-398981 fails in stress test",
|
||||
|
||||
@@ -39,33 +39,31 @@ 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
|
||||
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));
|
||||
void setNumBlocks(int blocksPerCU, int threadsPerBlock, size_t N, unsigned& blocks) {
|
||||
int device;
|
||||
HIP_CHECK_THREAD(hipGetDevice(&device));
|
||||
hipDeviceProp_t props;
|
||||
HIP_CHECK_THREAD(hipGetDeviceProperties(&props, device));
|
||||
|
||||
unsigned blocks = props.multiProcessorCount * blocksPerCU;
|
||||
if (blocks * threadsPerBlock > N) {
|
||||
blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
|
||||
}
|
||||
return blocks;
|
||||
blocks = props.multiProcessorCount * blocksPerCU;
|
||||
if (blocks * threadsPerBlock > N) {
|
||||
blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
|
||||
}
|
||||
}
|
||||
} // namespace MemcpyStream
|
||||
|
||||
|
||||
@@ -73,47 +71,40 @@ class HipMemcpyWithStreamMultiThreadtests {
|
||||
public:
|
||||
// Test hipMemcpyWithStream with one streams and launch kernel in
|
||||
// that stream, verify the data.
|
||||
void TestwithOnestream(bool &val_res);
|
||||
void TestwithOnestream(bool& val_res);
|
||||
// Test hipMemcpyWithStream with two streams and launch kernels in
|
||||
// two streams, verify the data.
|
||||
void TestwithTwoStream(bool &val_res);
|
||||
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);
|
||||
void TestOnMultiGPUwithOneStream(bool& val_res);
|
||||
// Test hipMemcpyWithStream to copy data from
|
||||
// device to host (hipMemcpyDeviceToHost).
|
||||
void TestkindDtoH(bool &val_res);
|
||||
void TestkindDtoH(bool& val_res);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on MultiGPU.
|
||||
void TestkindDtoD(bool &val_res);
|
||||
void TestkindDtoD(bool& val_res);
|
||||
// Test hipMemcpyWithStream with hipMemcpyHostToHost.
|
||||
void TestkindHtoH(bool &val_res);
|
||||
void TestkindHtoH(bool& val_res);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDefault.
|
||||
void TestkindDefault(bool &val_res);
|
||||
void TestkindDefault(bool& val_res);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDefault for
|
||||
// device to device transfer case.
|
||||
void TestkindDefaultForDtoD(bool &val_res);
|
||||
void TestkindDefaultForDtoD(bool& val_res);
|
||||
// Test hipMemcpyWithStream with hipMemcpyDeviceToDevice on same device.
|
||||
void TestDtoDonSameDevice(bool &val_res);
|
||||
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);
|
||||
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);
|
||||
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);
|
||||
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));
|
||||
void HipMemcpyWithStreamMultiThreadtests::AllocateMemory(int** A_d, int** B_d, int** C_d, int** A_h,
|
||||
int** B_h, int** C_h) {
|
||||
HIP_CHECK_THREAD(hipMalloc(A_d, Nbytes));
|
||||
HIP_CHECK_THREAD(hipMalloc(B_d, Nbytes));
|
||||
HIP_CHECK_THREAD(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));
|
||||
@@ -125,20 +116,18 @@ void HipMemcpyWithStreamMultiThreadtests::AllocateMemory(int** A_d, int** B_d,
|
||||
}
|
||||
}
|
||||
|
||||
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));
|
||||
void HipMemcpyWithStreamMultiThreadtests::DeAllocateMemory(int* A_d, int* B_d, int* C_d, int* A_h,
|
||||
int* B_h, int* C_h) {
|
||||
HIP_CHECK_THREAD(hipFree(A_d));
|
||||
HIP_CHECK_THREAD(hipFree(B_d));
|
||||
HIP_CHECK_THREAD(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 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]) {
|
||||
@@ -150,123 +139,113 @@ bool HipMemcpyWithStreamMultiThreadtests::
|
||||
}
|
||||
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestwithOnestream(bool &val_res) {
|
||||
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);
|
||||
unsigned blocks = 0;
|
||||
MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK_THREAD(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);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
HIP_CHECK_THREAD(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIP_CHECK_THREAD(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_THREAD(hipGetLastError());
|
||||
HIP_CHECK_THREAD(hipStreamSynchronize(stream));
|
||||
HIP_CHECK_THREAD(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));
|
||||
HIP_CHECK_THREAD(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestwithTwoStream(bool &val_res) {
|
||||
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);
|
||||
unsigned blocks = 0;
|
||||
MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
|
||||
|
||||
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]);
|
||||
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) {
|
||||
HIP_CHECK_THREAD(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) {
|
||||
HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes, hipMemcpyHostToDevice, stream[i]));
|
||||
HIP_CHECK_THREAD(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);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
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);
|
||||
HIP_CHECK_THREAD(hipGetLastError());
|
||||
}
|
||||
|
||||
for (int i=0; i < NoofStreams; ++i) {
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
for (int i = 0; i < NoofStreams; ++i) {
|
||||
HIP_CHECK_THREAD(hipStreamSynchronize(stream[i]));
|
||||
HIP_CHECK_THREAD(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) {
|
||||
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]));
|
||||
HIP_CHECK_THREAD(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestDtoDonSameDevice(bool &val_res) {
|
||||
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);
|
||||
unsigned blocks = 0;
|
||||
MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
|
||||
|
||||
AllocateMemory(&A_d[0], &B_d[0], &C_d[0],
|
||||
&A_h[0], &B_h[0], &C_h[0]);
|
||||
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]));
|
||||
for (int i = 0; i < NoofStreams; ++i) {
|
||||
HIP_CHECK_THREAD(hipSetDevice(0));
|
||||
HIP_CHECK_THREAD(hipStreamCreate(&stream[i]));
|
||||
}
|
||||
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipMalloc(&A_d[1], Nbytes));
|
||||
HIPCHECK(hipMalloc(&B_d[1], Nbytes));
|
||||
HIPCHECK(hipMalloc(&C_d[1], Nbytes));
|
||||
HIP_CHECK_THREAD(hipSetDevice(0));
|
||||
HIP_CHECK_THREAD(hipMalloc(&A_d[1], Nbytes));
|
||||
HIP_CHECK_THREAD(hipMalloc(&B_d[1], Nbytes));
|
||||
HIP_CHECK_THREAD(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]));
|
||||
HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
HIP_CHECK_THREAD(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]));
|
||||
HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[1], A_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[1]));
|
||||
HIP_CHECK_THREAD(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);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
for (int i = 0; i < NoofStreams; ++i) {
|
||||
HIP_CHECK_THREAD(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);
|
||||
HIP_CHECK_THREAD(hipGetLastError());
|
||||
}
|
||||
|
||||
for (int i=0; i < NoofStreams; ++i) {
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
for (int i = 0; i < NoofStreams; ++i) {
|
||||
HIP_CHECK_THREAD(hipSetDevice(0));
|
||||
HIP_CHECK_THREAD(hipStreamSynchronize(stream[i]));
|
||||
HIP_CHECK_THREAD(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
val_res = ValidateResult(A_h[0], B_h[0], C_h[i]);
|
||||
}
|
||||
|
||||
@@ -274,31 +253,31 @@ void HipMemcpyWithStreamMultiThreadtests::TestDtoDonSameDevice(bool &val_res) {
|
||||
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]));
|
||||
HIP_CHECK_THREAD(hipFree(A_d[1]));
|
||||
}
|
||||
if (B_d[1]) {
|
||||
HIPCHECK(hipFree(B_d[1]));
|
||||
HIP_CHECK_THREAD(hipFree(B_d[1]));
|
||||
}
|
||||
if (C_d[1]) {
|
||||
HIPCHECK(hipFree(C_d[1]));
|
||||
HIP_CHECK_THREAD(hipFree(C_d[1]));
|
||||
}
|
||||
if (C_h[1]) {
|
||||
free(C_h[1]);
|
||||
}
|
||||
|
||||
|
||||
for (int i=0; i < NoofStreams; ++i) {
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
for (int i = 0; i < NoofStreams; ++i) {
|
||||
HIP_CHECK_THREAD(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::
|
||||
TestOnMultiGPUwithOneStream(bool &val_res) {
|
||||
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));
|
||||
unsigned blocks = 0;
|
||||
MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
|
||||
HIP_CHECK_THREAD(hipGetDeviceCount(&numDevices));
|
||||
// If you have single GPU machine the return
|
||||
if (numDevices <= 1) {
|
||||
return;
|
||||
@@ -307,84 +286,78 @@ void HipMemcpyWithStreamMultiThreadtests::
|
||||
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) {
|
||||
HIP_CHECK_THREAD(hipSetDevice(i));
|
||||
HIP_CHECK_THREAD(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) {
|
||||
HIP_CHECK_THREAD(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) {
|
||||
HIP_CHECK_THREAD(hipSetDevice(i));
|
||||
HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[i], A_h[i], Nbytes, hipMemcpyHostToDevice, stream[i]));
|
||||
HIP_CHECK_THREAD(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);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
for (int i = 0; i < numDevices; ++i) {
|
||||
HIP_CHECK_THREAD(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);
|
||||
HIP_CHECK_THREAD(hipGetLastError());
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
for (int i = 0; i < numDevices; ++i) {
|
||||
HIP_CHECK_THREAD(hipSetDevice(i));
|
||||
HIP_CHECK_THREAD(hipStreamSynchronize(stream[i]));
|
||||
HIP_CHECK_THREAD(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));
|
||||
for (int i = 0; i < numDevices; ++i) {
|
||||
HIP_CHECK_THREAD(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]));
|
||||
HIP_CHECK_THREAD(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestkindDtoH(bool &val_res) {
|
||||
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);
|
||||
unsigned blocks = 0;
|
||||
MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
|
||||
AllocateMemory(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h);
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK_THREAD(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);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes,
|
||||
hipMemcpyDeviceToHost, stream));
|
||||
HIP_CHECK_THREAD(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream));
|
||||
HIP_CHECK_THREAD(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_THREAD(hipGetLastError());
|
||||
HIP_CHECK_THREAD(hipStreamSynchronize(stream));
|
||||
HIP_CHECK_THREAD(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));
|
||||
HIP_CHECK_THREAD(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestkindDtoD(bool &val_res) {
|
||||
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));
|
||||
unsigned blocks = 0;
|
||||
MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
|
||||
HIP_CHECK_THREAD(hipGetDeviceCount(&numDevices));
|
||||
// If you have single GPU machine the return
|
||||
if (numDevices <= 1) {
|
||||
return;
|
||||
@@ -394,116 +367,110 @@ void HipMemcpyWithStreamMultiThreadtests::TestkindDtoD(bool &val_res) {
|
||||
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) {
|
||||
HIP_CHECK_THREAD(hipSetDevice(i));
|
||||
HIP_CHECK_THREAD(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]);
|
||||
HIP_CHECK_THREAD(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));
|
||||
for (int i = 1; i < numDevices; ++i) {
|
||||
HIP_CHECK_THREAD(hipSetDevice(i))
|
||||
HIP_CHECK_THREAD(hipMalloc(&A_d[i], Nbytes));
|
||||
HIP_CHECK_THREAD(hipMalloc(&B_d[i], Nbytes));
|
||||
HIP_CHECK_THREAD(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]));
|
||||
HIP_CHECK_THREAD(hipSetDevice(0));
|
||||
HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
HIP_CHECK_THREAD(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]));
|
||||
for (int i = 1; i < numDevices; ++i) {
|
||||
HIP_CHECK_THREAD(hipSetDevice(i));
|
||||
HIP_CHECK_THREAD(
|
||||
hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, hipMemcpyDeviceToDevice, stream[i]));
|
||||
HIP_CHECK_THREAD(
|
||||
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);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
for (int i = 0; i < numDevices; ++i) {
|
||||
HIP_CHECK_THREAD(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);
|
||||
HIP_CHECK_THREAD(hipGetLastError());
|
||||
}
|
||||
|
||||
for (int i=0; i < numDevices; ++i) {
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
HIPCHECK(hipStreamSynchronize(stream[i]));
|
||||
HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
|
||||
for (int i = 0; i < numDevices; ++i) {
|
||||
HIP_CHECK_THREAD(hipSetDevice(i));
|
||||
HIP_CHECK_THREAD(hipStreamSynchronize(stream[i]));
|
||||
HIP_CHECK_THREAD(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]));
|
||||
HIP_CHECK_THREAD(hipStreamDestroy(stream[0]));
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
for (int i = 1; i < numDevices; ++i) {
|
||||
if (A_d[i]) {
|
||||
HIPCHECK(hipFree(A_d[i]));
|
||||
HIP_CHECK_THREAD(hipFree(A_d[i]));
|
||||
}
|
||||
if (B_d[i]) {
|
||||
HIPCHECK(hipFree(B_d[i]));
|
||||
HIP_CHECK_THREAD(hipFree(B_d[i]));
|
||||
}
|
||||
if (C_d[i]) {
|
||||
HIPCHECK(hipFree(C_d[i]));
|
||||
HIP_CHECK_THREAD(hipFree(C_d[i]));
|
||||
}
|
||||
if (C_h[i]) {
|
||||
free(C_h[i]);
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
HIP_CHECK_THREAD(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::
|
||||
TestkindDefault(bool &val_res) {
|
||||
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);
|
||||
unsigned blocks = 0;
|
||||
MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
|
||||
AllocateMemory(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h);
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK_THREAD(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);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIPCHECK(hipMemcpyWithStream(C_h, C_d, Nbytes, hipMemcpyDefault, stream));
|
||||
HIP_CHECK_THREAD(hipMemcpyWithStream(A_d, A_h, Nbytes, hipMemcpyDefault, stream));
|
||||
HIP_CHECK_THREAD(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_THREAD(hipGetLastError());
|
||||
HIP_CHECK_THREAD(hipStreamSynchronize(stream));
|
||||
HIP_CHECK_THREAD(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));
|
||||
HIP_CHECK_THREAD(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::
|
||||
TestkindDefaultForDtoD(bool &val_res) {
|
||||
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));
|
||||
unsigned blocks = 0;
|
||||
MemcpyStream::setNumBlocks(blocksPerCU, threadsPerBlock, N, blocks);
|
||||
HIP_CHECK_THREAD(hipGetDeviceCount(&numDevices));
|
||||
// Test case will not run on single GPU setup.
|
||||
if (numDevices <= 1) {
|
||||
return;
|
||||
@@ -513,76 +480,70 @@ void HipMemcpyWithStreamMultiThreadtests::
|
||||
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]);
|
||||
HIP_CHECK_THREAD(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));
|
||||
for (int i = 1; i < numDevices; ++i) {
|
||||
HIP_CHECK_THREAD(hipSetDevice(i));
|
||||
HIP_CHECK_THREAD(hipMalloc(&A_d[i], Nbytes));
|
||||
HIP_CHECK_THREAD(hipMalloc(&B_d[i], Nbytes));
|
||||
HIP_CHECK_THREAD(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]));
|
||||
for (int i = 0; i < numDevices; ++i) {
|
||||
HIP_CHECK_THREAD(hipSetDevice(i));
|
||||
HIP_CHECK_THREAD(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]));
|
||||
HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
|
||||
HIP_CHECK_THREAD(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 = 1; i < numDevices; ++i) {
|
||||
HIP_CHECK_THREAD(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, hipMemcpyDefault, stream[i]));
|
||||
HIP_CHECK_THREAD(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);
|
||||
HIP_CHECK(hipGetLastError());
|
||||
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);
|
||||
HIP_CHECK_THREAD(hipGetLastError());
|
||||
}
|
||||
|
||||
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));
|
||||
for (int i = 0; i < numDevices; ++i) {
|
||||
HIP_CHECK_THREAD(hipSetDevice(i)); // hipMemcpy will be on this device
|
||||
HIP_CHECK_THREAD(hipStreamSynchronize(stream[i]));
|
||||
HIP_CHECK_THREAD(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]));
|
||||
HIP_CHECK_THREAD(hipStreamDestroy(stream[0]));
|
||||
|
||||
for (int i=1; i < numDevices; ++i) {
|
||||
for (int i = 1; i < numDevices; ++i) {
|
||||
if (A_d[i]) {
|
||||
HIPCHECK(hipFree(A_d[i]));
|
||||
HIP_CHECK_THREAD(hipFree(A_d[i]));
|
||||
}
|
||||
if (B_d[i]) {
|
||||
HIPCHECK(hipFree(B_d[i]));
|
||||
HIP_CHECK_THREAD(hipFree(B_d[i]));
|
||||
}
|
||||
if (C_d[i]) {
|
||||
HIPCHECK(hipFree(C_d[i]));
|
||||
HIP_CHECK_THREAD(hipFree(C_d[i]));
|
||||
}
|
||||
if (C_h[i]) {
|
||||
free(C_h[i]);
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream[i]));
|
||||
HIP_CHECK_THREAD(hipStreamDestroy(stream[i]));
|
||||
}
|
||||
}
|
||||
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestkindHtoH(bool &val_res) {
|
||||
void HipMemcpyWithStreamMultiThreadtests::TestkindHtoH(bool& val_res) {
|
||||
size_t Nbytes = N * sizeof(int);
|
||||
int *A_h, *B_h;
|
||||
|
||||
@@ -597,15 +558,15 @@ void HipMemcpyWithStreamMultiThreadtests::TestkindHtoH(bool &val_res) {
|
||||
}
|
||||
|
||||
hipStream_t stream;
|
||||
HIPCHECK(hipStreamCreate(&stream));
|
||||
HIP_CHECK_THREAD(hipStreamCreate(&stream));
|
||||
|
||||
HIPCHECK(hipMemcpyWithStream(B_h, A_h, Nbytes, hipMemcpyHostToHost, stream));
|
||||
HIPCHECK(hipStreamSynchronize(stream));
|
||||
HIP_CHECK_THREAD(hipMemcpyWithStream(B_h, A_h, Nbytes, hipMemcpyHostToHost, stream));
|
||||
HIP_CHECK_THREAD(hipStreamSynchronize(stream));
|
||||
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
if ((A_h[i] != B_h[i])) {
|
||||
val_res = false;
|
||||
break;
|
||||
val_res = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -615,64 +576,56 @@ void HipMemcpyWithStreamMultiThreadtests::TestkindHtoH(bool &val_res) {
|
||||
if (B_h) {
|
||||
free(B_h);
|
||||
}
|
||||
HIPCHECK(hipStreamDestroy(stream));
|
||||
HIP_CHECK_THREAD(hipStreamDestroy(stream));
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMemcpyWithStream_MultiThread") {
|
||||
const auto Threadcount{10};
|
||||
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) {
|
||||
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) ) {
|
||||
switch (static_cast<ops>(op)) {
|
||||
case ops::TestwithOnestream:
|
||||
th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::
|
||||
TestwithOnestream,
|
||||
&tests, std::ref(ret_val[i]));
|
||||
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]));
|
||||
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]));
|
||||
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]));
|
||||
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]));
|
||||
th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestkindDtoD, &tests,
|
||||
std::ref(ret_val[i]));
|
||||
break;
|
||||
case ops::TestOnMultiGPUwithOneStream:
|
||||
th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::
|
||||
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]));
|
||||
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]));
|
||||
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]));
|
||||
th[i] = std::thread(&HipMemcpyWithStreamMultiThreadtests::TestDtoDonSameDevice, &tests,
|
||||
std::ref(ret_val[i]));
|
||||
break;
|
||||
default: {}
|
||||
default: {
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -680,6 +633,8 @@ TEST_CASE("Unit_hipMemcpyWithStream_MultiThread") {
|
||||
th[i].join();
|
||||
}
|
||||
|
||||
HIP_CHECK_THREAD_FINALIZE();
|
||||
|
||||
for (uint32_t i = 0; i < Threadcount; i++) {
|
||||
REQUIRE(ret_val[i] == true);
|
||||
}
|
||||
|
||||
Referens i nytt ärende
Block a user